Generated by Cython 3.0.2
Yellow lines hint at Python interaction.
Click on a line that starts with a "+" to see the C code that Cython generated for it.
Raw output: _le_roberts_cross_gradients.c
+001: # cython: infer_types=True, wraparound=False, nonecheck=False, boundscheck=False, cdivision=True, language_level=3, profile=False, autogen_pxd=False
__pyx_t_4 = __Pyx_PyDict_NewPresized(0); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 1, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); if (PyDict_SetItem(__pyx_d, __pyx_n_s_test, __pyx_t_4) < 0) __PYX_ERR(0, 1, __pyx_L1_error) __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0;
002:
+003: import numpy as np
__pyx_t_7 = __Pyx_ImportDottedModule(__pyx_n_s_numpy, NULL); if (unlikely(!__pyx_t_7)) __PYX_ERR(0, 3, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_7); if (PyDict_SetItem(__pyx_d, __pyx_n_s_np, __pyx_t_7) < 0) __PYX_ERR(0, 3, __pyx_L1_error) __Pyx_DECREF(__pyx_t_7); __pyx_t_7 = 0;
004: cimport numpy as np
+005: from .__opencl__ import cl, cl_array
__pyx_t_7 = PyList_New(2); if (unlikely(!__pyx_t_7)) __PYX_ERR(0, 5, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_7); __Pyx_INCREF(__pyx_n_s_cl); __Pyx_GIVEREF(__pyx_n_s_cl); if (__Pyx_PyList_SET_ITEM(__pyx_t_7, 0, __pyx_n_s_cl)) __PYX_ERR(0, 5, __pyx_L1_error); __Pyx_INCREF(__pyx_n_s_cl_array); __Pyx_GIVEREF(__pyx_n_s_cl_array); if (__Pyx_PyList_SET_ITEM(__pyx_t_7, 1, __pyx_n_s_cl_array)) __PYX_ERR(0, 5, __pyx_L1_error); __pyx_t_4 = __Pyx_Import(__pyx_n_s_opencl_2, __pyx_t_7, 1); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 5, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_DECREF(__pyx_t_7); __pyx_t_7 = 0; __pyx_t_7 = __Pyx_ImportFrom(__pyx_t_4, __pyx_n_s_cl); if (unlikely(!__pyx_t_7)) __PYX_ERR(0, 5, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_7); if (PyDict_SetItem(__pyx_d, __pyx_n_s_cl, __pyx_t_7) < 0) __PYX_ERR(0, 5, __pyx_L1_error) __Pyx_DECREF(__pyx_t_7); __pyx_t_7 = 0; __pyx_t_7 = __Pyx_ImportFrom(__pyx_t_4, __pyx_n_s_cl_array); if (unlikely(!__pyx_t_7)) __PYX_ERR(0, 5, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_7); if (PyDict_SetItem(__pyx_d, __pyx_n_s_cl_array, __pyx_t_7) < 0) __PYX_ERR(0, 5, __pyx_L1_error) __Pyx_DECREF(__pyx_t_7); __pyx_t_7 = 0; __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0;
+006: from .__liquid_engine__ import LiquidEngine
__pyx_t_4 = PyList_New(1); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 6, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_INCREF(__pyx_n_s_LiquidEngine); __Pyx_GIVEREF(__pyx_n_s_LiquidEngine); if (__Pyx_PyList_SET_ITEM(__pyx_t_4, 0, __pyx_n_s_LiquidEngine)) __PYX_ERR(0, 6, __pyx_L1_error); __pyx_t_7 = __Pyx_Import(__pyx_n_s_liquid_engine, __pyx_t_4, 1); if (unlikely(!__pyx_t_7)) __PYX_ERR(0, 6, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_7); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __pyx_t_4 = __Pyx_ImportFrom(__pyx_t_7, __pyx_n_s_LiquidEngine); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 6, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); if (PyDict_SetItem(__pyx_d, __pyx_n_s_LiquidEngine, __pyx_t_4) < 0) __PYX_ERR(0, 6, __pyx_L1_error) __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __Pyx_DECREF(__pyx_t_7); __pyx_t_7 = 0;
007: from cython.parallel import prange
+008: from .__interpolation_tools__ import check_image
__pyx_t_7 = PyList_New(1); if (unlikely(!__pyx_t_7)) __PYX_ERR(0, 8, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_7); __Pyx_INCREF(__pyx_n_s_check_image); __Pyx_GIVEREF(__pyx_n_s_check_image); if (__Pyx_PyList_SET_ITEM(__pyx_t_7, 0, __pyx_n_s_check_image)) __PYX_ERR(0, 8, __pyx_L1_error); __pyx_t_4 = __Pyx_Import(__pyx_n_s_interpolation_tools, __pyx_t_7, 1); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 8, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_DECREF(__pyx_t_7); __pyx_t_7 = 0; __pyx_t_7 = __Pyx_ImportFrom(__pyx_t_4, __pyx_n_s_check_image); if (unlikely(!__pyx_t_7)) __PYX_ERR(0, 8, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_7); if (PyDict_SetItem(__pyx_d, __pyx_n_s_check_image, __pyx_t_7) < 0) __PYX_ERR(0, 8, __pyx_L1_error) __Pyx_DECREF(__pyx_t_7); __pyx_t_7 = 0; __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0;
009:
010: cdef extern from "_c_gradients.h":
011: void _c_gradient_roberts_cross(float* pixels, float* GxArray, float* GyArray, int w, int h) nogil
012:
+013: class GradientRobertsCross(LiquidEngine):
__Pyx_GetModuleGlobalName(__pyx_t_4, __pyx_n_s_LiquidEngine); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 13, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __pyx_t_7 = PyTuple_New(1); if (unlikely(!__pyx_t_7)) __PYX_ERR(0, 13, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_7); __Pyx_GIVEREF(__pyx_t_4); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_7, 0, __pyx_t_4)) __PYX_ERR(0, 13, __pyx_L1_error); __pyx_t_4 = 0; __pyx_t_4 = __Pyx_PEP560_update_bases(__pyx_t_7); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 13, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __pyx_t_5 = __Pyx_CalculateMetaclass(NULL, __pyx_t_4); if (unlikely(!__pyx_t_5)) __PYX_ERR(0, 13, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_5); __pyx_t_9 = __Pyx_Py3MetaclassPrepare(__pyx_t_5, __pyx_t_4, __pyx_n_s_GradientRobertsCross, __pyx_n_s_GradientRobertsCross, (PyObject *) NULL, __pyx_n_s_nanopyx_liquid__le_roberts_cross, (PyObject *) NULL); if (unlikely(!__pyx_t_9)) __PYX_ERR(0, 13, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_9); if (__pyx_t_4 != __pyx_t_7) { if (unlikely((PyDict_SetItemString(__pyx_t_9, "__orig_bases__", __pyx_t_7) < 0))) __PYX_ERR(0, 13, __pyx_L1_error) } __Pyx_DECREF(__pyx_t_7); __pyx_t_7 = 0; __pyx_t_7 = PyList_New(0); if (unlikely(!__pyx_t_7)) __PYX_ERR(0, 13, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_7); /* … */ __pyx_t_10 = __Pyx_Py3ClassCreate(__pyx_t_5, __pyx_n_s_GradientRobertsCross, __pyx_t_4, __pyx_t_9, NULL, 0, 0); if (unlikely(!__pyx_t_10)) __PYX_ERR(0, 13, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_10); if (__Pyx_CyFunction_InitClassCell(__pyx_t_7, __pyx_t_10) < 0) __PYX_ERR(0, 13, __pyx_L1_error) __Pyx_DECREF(__pyx_t_7); __pyx_t_7 = 0; if (PyDict_SetItem(__pyx_d, __pyx_n_s_GradientRobertsCross, __pyx_t_10) < 0) __PYX_ERR(0, 13, __pyx_L1_error) __Pyx_DECREF(__pyx_t_10); __pyx_t_10 = 0; __Pyx_DECREF(__pyx_t_9); __pyx_t_9 = 0; __Pyx_DECREF(__pyx_t_5); __pyx_t_5 = 0; __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0;
014:
+015: def __init__(self, clear_benchmarks=False, testing=False):
/* Python wrapper */
static PyObject *__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_1__init__(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
); /*proto*/
static PyMethodDef __pyx_mdef_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_1__init__ = {"__init__", (PyCFunction)(void*)(__Pyx_PyCFunction_FastCallWithKeywords)__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_1__init__, __Pyx_METH_FASTCALL|METH_KEYWORDS, 0};
static PyObject *__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_1__init__(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
) {
PyObject *__pyx_v_self = 0;
PyObject *__pyx_v_clear_benchmarks = 0;
PyObject *__pyx_v_testing = 0;
#if !CYTHON_METH_FASTCALL
CYTHON_UNUSED Py_ssize_t __pyx_nargs;
#endif
CYTHON_UNUSED PyObject *const *__pyx_kwvalues;
PyObject *__pyx_r = 0;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("__init__ (wrapper)", 0);
#if !CYTHON_METH_FASTCALL
#if CYTHON_ASSUME_SAFE_MACROS
__pyx_nargs = PyTuple_GET_SIZE(__pyx_args);
#else
__pyx_nargs = PyTuple_Size(__pyx_args);
if (unlikely((__pyx_nargs < 0))) __PYX_ERR(0, 15, __pyx_L3_error)
#endif
#endif
__pyx_kwvalues = __Pyx_KwValues_FASTCALL(__pyx_args, __pyx_nargs);
{
PyObject **__pyx_pyargnames[] = {&__pyx_n_s_self,&__pyx_n_s_clear_benchmarks,&__pyx_n_s_testing,0};
PyObject* values[3] = {0,0,0};
values[1] = __Pyx_Arg_NewRef_FASTCALL(((PyObject *)((PyObject *)Py_False)));
values[2] = __Pyx_Arg_NewRef_FASTCALL(((PyObject *)((PyObject *)Py_False)));
if (__pyx_kwds) {
Py_ssize_t kw_args;
switch (__pyx_nargs) {
case 3: values[2] = __Pyx_Arg_FASTCALL(__pyx_args, 2);
CYTHON_FALLTHROUGH;
case 2: values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
CYTHON_FALLTHROUGH;
case 1: values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
CYTHON_FALLTHROUGH;
case 0: break;
default: goto __pyx_L5_argtuple_error;
}
kw_args = __Pyx_NumKwargs_FASTCALL(__pyx_kwds);
switch (__pyx_nargs) {
case 0:
if (likely((values[0] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_self)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[0]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 15, __pyx_L3_error)
else goto __pyx_L5_argtuple_error;
CYTHON_FALLTHROUGH;
case 1:
if (kw_args > 0) {
PyObject* value = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_clear_benchmarks);
if (value) { values[1] = __Pyx_Arg_NewRef_FASTCALL(value); kw_args--; }
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 15, __pyx_L3_error)
}
CYTHON_FALLTHROUGH;
case 2:
if (kw_args > 0) {
PyObject* value = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_testing);
if (value) { values[2] = __Pyx_Arg_NewRef_FASTCALL(value); kw_args--; }
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 15, __pyx_L3_error)
}
}
if (unlikely(kw_args > 0)) {
const Py_ssize_t kwd_pos_args = __pyx_nargs;
if (unlikely(__Pyx_ParseOptionalKeywords(__pyx_kwds, __pyx_kwvalues, __pyx_pyargnames, 0, values + 0, kwd_pos_args, "__init__") < 0)) __PYX_ERR(0, 15, __pyx_L3_error)
}
} else {
switch (__pyx_nargs) {
case 3: values[2] = __Pyx_Arg_FASTCALL(__pyx_args, 2);
CYTHON_FALLTHROUGH;
case 2: values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
CYTHON_FALLTHROUGH;
case 1: values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
break;
default: goto __pyx_L5_argtuple_error;
}
}
__pyx_v_self = values[0];
__pyx_v_clear_benchmarks = values[1];
__pyx_v_testing = values[2];
}
goto __pyx_L4_argument_unpacking_done;
__pyx_L5_argtuple_error:;
__Pyx_RaiseArgtupleInvalid("__init__", 0, 1, 3, __pyx_nargs); __PYX_ERR(0, 15, __pyx_L3_error)
goto __pyx_L3_error;
__pyx_L3_error:;
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__Pyx_AddTraceback("nanopyx.liquid._le_roberts_cross_gradients.GradientRobertsCross.__init__", __pyx_clineno, __pyx_lineno, __pyx_filename);
__Pyx_RefNannyFinishContext();
return NULL;
__pyx_L4_argument_unpacking_done:;
__pyx_r = __pyx_pf_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross___init__(__pyx_self, __pyx_v_self, __pyx_v_clear_benchmarks, __pyx_v_testing);
/* function exit code */
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
static PyObject *__pyx_pf_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross___init__(CYTHON_UNUSED PyObject *__pyx_self, PyObject *__pyx_v_self, PyObject *__pyx_v_clear_benchmarks, PyObject *__pyx_v_testing) {
PyObject *__pyx_r = NULL;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("__init__", 0);
/* … */
/* function exit code */
__pyx_r = Py_None; __Pyx_INCREF(Py_None);
goto __pyx_L0;
__pyx_L1_error:;
__Pyx_XDECREF(__pyx_t_1);
__Pyx_XDECREF(__pyx_t_2);
__Pyx_XDECREF(__pyx_t_3);
__Pyx_AddTraceback("nanopyx.liquid._le_roberts_cross_gradients.GradientRobertsCross.__init__", __pyx_clineno, __pyx_lineno, __pyx_filename);
__pyx_r = NULL;
__pyx_L0:;
__Pyx_XGIVEREF(__pyx_r);
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
/* … */
__pyx_tuple__22 = PyTuple_Pack(3, __pyx_n_s_self, __pyx_n_s_clear_benchmarks, __pyx_n_s_testing); if (unlikely(!__pyx_tuple__22)) __PYX_ERR(0, 15, __pyx_L1_error)
__Pyx_GOTREF(__pyx_tuple__22);
__Pyx_GIVEREF(__pyx_tuple__22);
__pyx_codeobj__23 = (PyObject*)__Pyx_PyCode_New(3, 0, 0, 3, 0, CO_OPTIMIZED|CO_NEWLOCALS, __pyx_empty_bytes, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_tuple__22, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_kp_s_src_nanopyx_liquid__le_roberts_c, __pyx_n_s_init, 15, __pyx_empty_bytes); if (unlikely(!__pyx_codeobj__23)) __PYX_ERR(0, 15, __pyx_L1_error)
/* … */
__pyx_t_10 = __Pyx_CyFunction_New(&__pyx_mdef_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_1__init__, 0, __pyx_n_s_GradientRobertsCross___init, NULL, __pyx_n_s_nanopyx_liquid__le_roberts_cross, __pyx_d, ((PyObject *)__pyx_codeobj__23)); if (unlikely(!__pyx_t_10)) __PYX_ERR(0, 15, __pyx_L1_error)
__Pyx_GOTREF(__pyx_t_10);
__Pyx_INCREF(__pyx_t_10);
PyList_Append(__pyx_t_7, __pyx_t_10);
__Pyx_GIVEREF(__pyx_t_10);
__Pyx_CyFunction_SetDefaultsTuple(__pyx_t_10, __pyx_tuple__24);
if (__Pyx_SetNameInClass(__pyx_t_9, __pyx_n_s_init, __pyx_t_10) < 0) __PYX_ERR(0, 15, __pyx_L1_error)
__Pyx_DECREF(__pyx_t_10); __pyx_t_10 = 0;
__pyx_tuple__24 = PyTuple_Pack(2, ((PyObject *)Py_False), ((PyObject *)Py_False)); if (unlikely(!__pyx_tuple__24)) __PYX_ERR(0, 15, __pyx_L1_error)
__Pyx_GOTREF(__pyx_tuple__24);
__Pyx_GIVEREF(__pyx_tuple__24);
+016: self._designation = "GradientRobertsCross"
if (__Pyx_PyObject_SetAttrStr(__pyx_v_self, __pyx_n_s_designation, __pyx_n_u_GradientRobertsCross) < 0) __PYX_ERR(0, 16, __pyx_L1_error)
+017: super().__init__(clear_benchmarks=clear_benchmarks, testing=testing,
__pyx_t_1 = __Pyx_CyFunction_GetClassObj(__pyx_self); if (!__pyx_t_1) { PyErr_SetString(PyExc_SystemError, "super(): empty __class__ cell"); __PYX_ERR(0, 17, __pyx_L1_error) } __Pyx_INCREF(__pyx_t_1); __pyx_t_2 = PyTuple_New(2); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 17, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_GIVEREF(__pyx_t_1); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_2, 0, __pyx_t_1)) __PYX_ERR(0, 17, __pyx_L1_error); __Pyx_INCREF(__pyx_v_self); __Pyx_GIVEREF(__pyx_v_self); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_2, 1, __pyx_v_self)) __PYX_ERR(0, 17, __pyx_L1_error); __pyx_t_1 = 0; __pyx_t_1 = __Pyx_PyObject_Call(__pyx_builtin_super, __pyx_t_2, NULL); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 17, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = __Pyx_PyObject_GetAttrStr(__pyx_t_1, __pyx_n_s_init); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 17, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_t_1 = __Pyx_PyDict_NewPresized(8); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 17, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); if (PyDict_SetItem(__pyx_t_1, __pyx_n_s_clear_benchmarks, __pyx_v_clear_benchmarks) < 0) __PYX_ERR(0, 17, __pyx_L1_error) if (PyDict_SetItem(__pyx_t_1, __pyx_n_s_testing, __pyx_v_testing) < 0) __PYX_ERR(0, 17, __pyx_L1_error) /* … */ __pyx_t_3 = __Pyx_PyObject_Call(__pyx_t_2, __pyx_empty_tuple, __pyx_t_1); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 17, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0;
+018: unthreaded_=True, threaded_=True, threaded_static_=True,
if (PyDict_SetItem(__pyx_t_1, __pyx_n_s_unthreaded, Py_True) < 0) __PYX_ERR(0, 17, __pyx_L1_error) if (PyDict_SetItem(__pyx_t_1, __pyx_n_s_threaded, Py_True) < 0) __PYX_ERR(0, 17, __pyx_L1_error) if (PyDict_SetItem(__pyx_t_1, __pyx_n_s_threaded_static, Py_True) < 0) __PYX_ERR(0, 17, __pyx_L1_error)
+019: threaded_dynamic_=True, threaded_guided_=True, opencl_=True)
if (PyDict_SetItem(__pyx_t_1, __pyx_n_s_threaded_dynamic, Py_True) < 0) __PYX_ERR(0, 17, __pyx_L1_error) if (PyDict_SetItem(__pyx_t_1, __pyx_n_s_threaded_guided, Py_True) < 0) __PYX_ERR(0, 17, __pyx_L1_error) if (PyDict_SetItem(__pyx_t_1, __pyx_n_s_opencl, Py_True) < 0) __PYX_ERR(0, 17, __pyx_L1_error)
020:
+021: self._default_benchmarks = {'OpenCL': {"(['shape(100, 100, 100)'], {})": [1000000, 0.09102950000033161, 0.013413084000148956, 0.013892708000184939], "(['shape(500, 500, 500)'], {})": [125000000, 0.4193332919999193, 0.36434012500012614, 0.3667828749998989]}, 'Threaded': {"(['shape(100, 100, 100)'], {})": [1000000, 0.001157208999757131, 0.0011499580000418064, 0.001164416999927198], "(['shape(500, 500, 500)'], {})": [125000000, 0.1371792079999068, 0.12759208300030878, 0.1256785000000491]}, 'Threaded_dynamic': {"(['shape(100, 100, 100)'], {})": [1000000, 0.0009977079998861882, 0.0009290419998251309, 0.0009676250001575681], "(['shape(500, 500, 500)'], {})": [125000000, 0.11562845800017385, 0.12086337500022637, 0.11272079200034568]}, 'Threaded_guided': {"(['shape(100, 100, 100)'], {})": [1000000, 0.0009116250002989545, 0.0009819160000006377, 0.000968583000030776], "(['shape(500, 500, 500)'], {})": [125000000, 0.11294766700029868, 0.11097495800004253, 0.10783799999990151]}, 'Threaded_static': {"(['shape(100, 100, 100)'], {})": [1000000, 0.000992541999949026, 0.0010514160003367579, 0.0011669999998957792], "(['shape(500, 500, 500)'], {})": [125000000, 0.12436966700033736, 0.11687999999958265, 0.10908783299964853]}, 'Unthreaded': {"(['shape(100, 100, 100)'], {})": [1000000, 0.0016904580002119474, 0.001859041999978217, 0.0016570830002820003], "(['shape(500, 500, 500)'], {})": [125000000, 0.22857279200025005, 0.22197804200004612, 0.21746020799992039]}}
__pyx_t_3 = __Pyx_PyDict_NewPresized(6); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_1 = __Pyx_PyDict_NewPresized(2); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_2 = PyList_New(4); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_int_1000000); __Pyx_GIVEREF(__pyx_int_1000000); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 0, __pyx_int_1000000)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_09102950000033161); __Pyx_GIVEREF(__pyx_float_0_09102950000033161); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 1, __pyx_float_0_09102950000033161)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_013413084000148956); __Pyx_GIVEREF(__pyx_float_0_013413084000148956); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 2, __pyx_float_0_013413084000148956)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_013892708000184939); __Pyx_GIVEREF(__pyx_float_0_013892708000184939); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 3, __pyx_float_0_013892708000184939)) __PYX_ERR(0, 21, __pyx_L1_error); if (PyDict_SetItem(__pyx_t_1, __pyx_kp_u_shape_100_100_100, __pyx_t_2) < 0) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = PyList_New(4); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_int_125000000); __Pyx_GIVEREF(__pyx_int_125000000); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 0, __pyx_int_125000000)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_4193332919999193); __Pyx_GIVEREF(__pyx_float_0_4193332919999193); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 1, __pyx_float_0_4193332919999193)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_36434012500012614); __Pyx_GIVEREF(__pyx_float_0_36434012500012614); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 2, __pyx_float_0_36434012500012614)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_3667828749998989); __Pyx_GIVEREF(__pyx_float_0_3667828749998989); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 3, __pyx_float_0_3667828749998989)) __PYX_ERR(0, 21, __pyx_L1_error); if (PyDict_SetItem(__pyx_t_1, __pyx_kp_u_shape_500_500_500, __pyx_t_2) < 0) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; if (PyDict_SetItem(__pyx_t_3, __pyx_n_u_OpenCL, __pyx_t_1) < 0) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_t_1 = __Pyx_PyDict_NewPresized(2); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_2 = PyList_New(4); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_int_1000000); __Pyx_GIVEREF(__pyx_int_1000000); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 0, __pyx_int_1000000)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_001157208999757131); __Pyx_GIVEREF(__pyx_float_0_001157208999757131); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 1, __pyx_float_0_001157208999757131)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_0011499580000418064); __Pyx_GIVEREF(__pyx_float_0_0011499580000418064); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 2, __pyx_float_0_0011499580000418064)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_001164416999927198); __Pyx_GIVEREF(__pyx_float_0_001164416999927198); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 3, __pyx_float_0_001164416999927198)) __PYX_ERR(0, 21, __pyx_L1_error); if (PyDict_SetItem(__pyx_t_1, __pyx_kp_u_shape_100_100_100, __pyx_t_2) < 0) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = PyList_New(4); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_int_125000000); __Pyx_GIVEREF(__pyx_int_125000000); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 0, __pyx_int_125000000)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_1371792079999068); __Pyx_GIVEREF(__pyx_float_0_1371792079999068); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 1, __pyx_float_0_1371792079999068)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_12759208300030878); __Pyx_GIVEREF(__pyx_float_0_12759208300030878); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 2, __pyx_float_0_12759208300030878)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_1256785000000491); __Pyx_GIVEREF(__pyx_float_0_1256785000000491); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 3, __pyx_float_0_1256785000000491)) __PYX_ERR(0, 21, __pyx_L1_error); if (PyDict_SetItem(__pyx_t_1, __pyx_kp_u_shape_500_500_500, __pyx_t_2) < 0) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; if (PyDict_SetItem(__pyx_t_3, __pyx_n_u_Threaded, __pyx_t_1) < 0) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_t_1 = __Pyx_PyDict_NewPresized(2); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_2 = PyList_New(4); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_int_1000000); __Pyx_GIVEREF(__pyx_int_1000000); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 0, __pyx_int_1000000)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_0009977079998861882); __Pyx_GIVEREF(__pyx_float_0_0009977079998861882); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 1, __pyx_float_0_0009977079998861882)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_0009290419998251309); __Pyx_GIVEREF(__pyx_float_0_0009290419998251309); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 2, __pyx_float_0_0009290419998251309)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_0009676250001575681); __Pyx_GIVEREF(__pyx_float_0_0009676250001575681); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 3, __pyx_float_0_0009676250001575681)) __PYX_ERR(0, 21, __pyx_L1_error); if (PyDict_SetItem(__pyx_t_1, __pyx_kp_u_shape_100_100_100, __pyx_t_2) < 0) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = PyList_New(4); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_int_125000000); __Pyx_GIVEREF(__pyx_int_125000000); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 0, __pyx_int_125000000)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_11562845800017385); __Pyx_GIVEREF(__pyx_float_0_11562845800017385); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 1, __pyx_float_0_11562845800017385)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_12086337500022637); __Pyx_GIVEREF(__pyx_float_0_12086337500022637); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 2, __pyx_float_0_12086337500022637)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_11272079200034568); __Pyx_GIVEREF(__pyx_float_0_11272079200034568); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 3, __pyx_float_0_11272079200034568)) __PYX_ERR(0, 21, __pyx_L1_error); if (PyDict_SetItem(__pyx_t_1, __pyx_kp_u_shape_500_500_500, __pyx_t_2) < 0) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; if (PyDict_SetItem(__pyx_t_3, __pyx_n_u_Threaded_dynamic, __pyx_t_1) < 0) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_t_1 = __Pyx_PyDict_NewPresized(2); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_2 = PyList_New(4); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_int_1000000); __Pyx_GIVEREF(__pyx_int_1000000); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 0, __pyx_int_1000000)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_0009116250002989545); __Pyx_GIVEREF(__pyx_float_0_0009116250002989545); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 1, __pyx_float_0_0009116250002989545)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_0009819160000006377); __Pyx_GIVEREF(__pyx_float_0_0009819160000006377); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 2, __pyx_float_0_0009819160000006377)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_000968583000030776); __Pyx_GIVEREF(__pyx_float_0_000968583000030776); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 3, __pyx_float_0_000968583000030776)) __PYX_ERR(0, 21, __pyx_L1_error); if (PyDict_SetItem(__pyx_t_1, __pyx_kp_u_shape_100_100_100, __pyx_t_2) < 0) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = PyList_New(4); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_int_125000000); __Pyx_GIVEREF(__pyx_int_125000000); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 0, __pyx_int_125000000)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_11294766700029868); __Pyx_GIVEREF(__pyx_float_0_11294766700029868); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 1, __pyx_float_0_11294766700029868)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_11097495800004253); __Pyx_GIVEREF(__pyx_float_0_11097495800004253); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 2, __pyx_float_0_11097495800004253)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_10783799999990151); __Pyx_GIVEREF(__pyx_float_0_10783799999990151); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 3, __pyx_float_0_10783799999990151)) __PYX_ERR(0, 21, __pyx_L1_error); if (PyDict_SetItem(__pyx_t_1, __pyx_kp_u_shape_500_500_500, __pyx_t_2) < 0) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; if (PyDict_SetItem(__pyx_t_3, __pyx_n_u_Threaded_guided, __pyx_t_1) < 0) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_t_1 = __Pyx_PyDict_NewPresized(2); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_2 = PyList_New(4); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_int_1000000); __Pyx_GIVEREF(__pyx_int_1000000); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 0, __pyx_int_1000000)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_000992541999949026); __Pyx_GIVEREF(__pyx_float_0_000992541999949026); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 1, __pyx_float_0_000992541999949026)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_0010514160003367579); __Pyx_GIVEREF(__pyx_float_0_0010514160003367579); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 2, __pyx_float_0_0010514160003367579)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_0011669999998957792); __Pyx_GIVEREF(__pyx_float_0_0011669999998957792); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 3, __pyx_float_0_0011669999998957792)) __PYX_ERR(0, 21, __pyx_L1_error); if (PyDict_SetItem(__pyx_t_1, __pyx_kp_u_shape_100_100_100, __pyx_t_2) < 0) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = PyList_New(4); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_int_125000000); __Pyx_GIVEREF(__pyx_int_125000000); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 0, __pyx_int_125000000)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_12436966700033736); __Pyx_GIVEREF(__pyx_float_0_12436966700033736); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 1, __pyx_float_0_12436966700033736)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_11687999999958265); __Pyx_GIVEREF(__pyx_float_0_11687999999958265); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 2, __pyx_float_0_11687999999958265)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_10908783299964853); __Pyx_GIVEREF(__pyx_float_0_10908783299964853); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 3, __pyx_float_0_10908783299964853)) __PYX_ERR(0, 21, __pyx_L1_error); if (PyDict_SetItem(__pyx_t_1, __pyx_kp_u_shape_500_500_500, __pyx_t_2) < 0) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; if (PyDict_SetItem(__pyx_t_3, __pyx_n_u_Threaded_static, __pyx_t_1) < 0) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_t_1 = __Pyx_PyDict_NewPresized(2); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_2 = PyList_New(4); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_int_1000000); __Pyx_GIVEREF(__pyx_int_1000000); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 0, __pyx_int_1000000)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_0016904580002119474); __Pyx_GIVEREF(__pyx_float_0_0016904580002119474); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 1, __pyx_float_0_0016904580002119474)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_001859041999978217); __Pyx_GIVEREF(__pyx_float_0_001859041999978217); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 2, __pyx_float_0_001859041999978217)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_0016570830002820003); __Pyx_GIVEREF(__pyx_float_0_0016570830002820003); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 3, __pyx_float_0_0016570830002820003)) __PYX_ERR(0, 21, __pyx_L1_error); if (PyDict_SetItem(__pyx_t_1, __pyx_kp_u_shape_100_100_100, __pyx_t_2) < 0) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = PyList_New(4); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_int_125000000); __Pyx_GIVEREF(__pyx_int_125000000); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 0, __pyx_int_125000000)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_22857279200025005); __Pyx_GIVEREF(__pyx_float_0_22857279200025005); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 1, __pyx_float_0_22857279200025005)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_22197804200004612); __Pyx_GIVEREF(__pyx_float_0_22197804200004612); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 2, __pyx_float_0_22197804200004612)) __PYX_ERR(0, 21, __pyx_L1_error); __Pyx_INCREF(__pyx_float_0_21746020799992039); __Pyx_GIVEREF(__pyx_float_0_21746020799992039); if (__Pyx_PyList_SET_ITEM(__pyx_t_2, 3, __pyx_float_0_21746020799992039)) __PYX_ERR(0, 21, __pyx_L1_error); if (PyDict_SetItem(__pyx_t_1, __pyx_kp_u_shape_500_500_500, __pyx_t_2) < 0) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; if (PyDict_SetItem(__pyx_t_3, __pyx_n_u_Unthreaded, __pyx_t_1) < 0) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; if (__Pyx_PyObject_SetAttrStr(__pyx_v_self, __pyx_n_s_default_benchmarks, __pyx_t_3) < 0) __PYX_ERR(0, 21, __pyx_L1_error) __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0;
022:
+023: def run(self, image, run_type = None):
/* Python wrapper */
static PyObject *__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_3run(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
); /*proto*/
static PyMethodDef __pyx_mdef_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_3run = {"run", (PyCFunction)(void*)(__Pyx_PyCFunction_FastCallWithKeywords)__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_3run, __Pyx_METH_FASTCALL|METH_KEYWORDS, 0};
static PyObject *__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_3run(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
) {
PyObject *__pyx_v_self = 0;
PyObject *__pyx_v_image = 0;
PyObject *__pyx_v_run_type = 0;
#if !CYTHON_METH_FASTCALL
CYTHON_UNUSED Py_ssize_t __pyx_nargs;
#endif
CYTHON_UNUSED PyObject *const *__pyx_kwvalues;
PyObject *__pyx_r = 0;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("run (wrapper)", 0);
#if !CYTHON_METH_FASTCALL
#if CYTHON_ASSUME_SAFE_MACROS
__pyx_nargs = PyTuple_GET_SIZE(__pyx_args);
#else
__pyx_nargs = PyTuple_Size(__pyx_args);
if (unlikely((__pyx_nargs < 0))) __PYX_ERR(0, 23, __pyx_L3_error)
#endif
#endif
__pyx_kwvalues = __Pyx_KwValues_FASTCALL(__pyx_args, __pyx_nargs);
{
PyObject **__pyx_pyargnames[] = {&__pyx_n_s_self,&__pyx_n_s_image,&__pyx_n_s_run_type,0};
PyObject* values[3] = {0,0,0};
values[2] = __Pyx_Arg_NewRef_FASTCALL(((PyObject *)Py_None));
if (__pyx_kwds) {
Py_ssize_t kw_args;
switch (__pyx_nargs) {
case 3: values[2] = __Pyx_Arg_FASTCALL(__pyx_args, 2);
CYTHON_FALLTHROUGH;
case 2: values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
CYTHON_FALLTHROUGH;
case 1: values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
CYTHON_FALLTHROUGH;
case 0: break;
default: goto __pyx_L5_argtuple_error;
}
kw_args = __Pyx_NumKwargs_FASTCALL(__pyx_kwds);
switch (__pyx_nargs) {
case 0:
if (likely((values[0] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_self)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[0]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 23, __pyx_L3_error)
else goto __pyx_L5_argtuple_error;
CYTHON_FALLTHROUGH;
case 1:
if (likely((values[1] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_image)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[1]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 23, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("run", 0, 2, 3, 1); __PYX_ERR(0, 23, __pyx_L3_error)
}
CYTHON_FALLTHROUGH;
case 2:
if (kw_args > 0) {
PyObject* value = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_run_type);
if (value) { values[2] = __Pyx_Arg_NewRef_FASTCALL(value); kw_args--; }
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 23, __pyx_L3_error)
}
}
if (unlikely(kw_args > 0)) {
const Py_ssize_t kwd_pos_args = __pyx_nargs;
if (unlikely(__Pyx_ParseOptionalKeywords(__pyx_kwds, __pyx_kwvalues, __pyx_pyargnames, 0, values + 0, kwd_pos_args, "run") < 0)) __PYX_ERR(0, 23, __pyx_L3_error)
}
} else {
switch (__pyx_nargs) {
case 3: values[2] = __Pyx_Arg_FASTCALL(__pyx_args, 2);
CYTHON_FALLTHROUGH;
case 2: values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
break;
default: goto __pyx_L5_argtuple_error;
}
}
__pyx_v_self = values[0];
__pyx_v_image = values[1];
__pyx_v_run_type = values[2];
}
goto __pyx_L4_argument_unpacking_done;
__pyx_L5_argtuple_error:;
__Pyx_RaiseArgtupleInvalid("run", 0, 2, 3, __pyx_nargs); __PYX_ERR(0, 23, __pyx_L3_error)
goto __pyx_L3_error;
__pyx_L3_error:;
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__Pyx_AddTraceback("nanopyx.liquid._le_roberts_cross_gradients.GradientRobertsCross.run", __pyx_clineno, __pyx_lineno, __pyx_filename);
__Pyx_RefNannyFinishContext();
return NULL;
__pyx_L4_argument_unpacking_done:;
__pyx_r = __pyx_pf_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_2run(__pyx_self, __pyx_v_self, __pyx_v_image, __pyx_v_run_type);
/* function exit code */
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
static PyObject *__pyx_pf_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_2run(CYTHON_UNUSED PyObject *__pyx_self, PyObject *__pyx_v_self, PyObject *__pyx_v_image, PyObject *__pyx_v_run_type) {
PyObject *__pyx_r = NULL;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("run", 0);
__Pyx_INCREF(__pyx_v_image);
/* … */
/* function exit code */
__pyx_L1_error:;
__Pyx_XDECREF(__pyx_t_1);
__Pyx_XDECREF(__pyx_t_2);
__Pyx_XDECREF(__pyx_t_3);
__Pyx_XDECREF(__pyx_t_5);
__Pyx_AddTraceback("nanopyx.liquid._le_roberts_cross_gradients.GradientRobertsCross.run", __pyx_clineno, __pyx_lineno, __pyx_filename);
__pyx_r = NULL;
__pyx_L0:;
__Pyx_XDECREF(__pyx_v_image);
__Pyx_XGIVEREF(__pyx_r);
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
/* … */
__pyx_tuple__25 = PyTuple_Pack(3, __pyx_n_s_self, __pyx_n_s_image, __pyx_n_s_run_type); if (unlikely(!__pyx_tuple__25)) __PYX_ERR(0, 23, __pyx_L1_error)
__Pyx_GOTREF(__pyx_tuple__25);
__Pyx_GIVEREF(__pyx_tuple__25);
__pyx_codeobj__26 = (PyObject*)__Pyx_PyCode_New(3, 0, 0, 3, 0, CO_OPTIMIZED|CO_NEWLOCALS, __pyx_empty_bytes, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_tuple__25, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_kp_s_src_nanopyx_liquid__le_roberts_c, __pyx_n_s_run_2, 23, __pyx_empty_bytes); if (unlikely(!__pyx_codeobj__26)) __PYX_ERR(0, 23, __pyx_L1_error)
/* … */
__pyx_t_10 = __Pyx_CyFunction_New(&__pyx_mdef_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_3run, 0, __pyx_n_s_GradientRobertsCross_run, NULL, __pyx_n_s_nanopyx_liquid__le_roberts_cross, __pyx_d, ((PyObject *)__pyx_codeobj__26)); if (unlikely(!__pyx_t_10)) __PYX_ERR(0, 23, __pyx_L1_error)
__Pyx_GOTREF(__pyx_t_10);
__Pyx_CyFunction_SetDefaultsTuple(__pyx_t_10, __pyx_tuple__27);
if (__Pyx_SetNameInClass(__pyx_t_9, __pyx_n_s_run_2, __pyx_t_10) < 0) __PYX_ERR(0, 23, __pyx_L1_error)
__Pyx_DECREF(__pyx_t_10); __pyx_t_10 = 0;
__pyx_tuple__27 = PyTuple_Pack(1, Py_None); if (unlikely(!__pyx_tuple__27)) __PYX_ERR(0, 23, __pyx_L1_error)
__Pyx_GOTREF(__pyx_tuple__27);
__Pyx_GIVEREF(__pyx_tuple__27);
+024: image = check_image(image)
__Pyx_GetModuleGlobalName(__pyx_t_2, __pyx_n_s_check_image); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 24, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_t_3 = NULL; __pyx_t_4 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_2))) { __pyx_t_3 = PyMethod_GET_SELF(__pyx_t_2); if (likely(__pyx_t_3)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_2); __Pyx_INCREF(__pyx_t_3); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_2, function); __pyx_t_4 = 1; } } #endif { PyObject *__pyx_callargs[2] = {__pyx_t_3, __pyx_v_image}; __pyx_t_1 = __Pyx_PyObject_FastCall(__pyx_t_2, __pyx_callargs+1-__pyx_t_4, 1+__pyx_t_4); __Pyx_XDECREF(__pyx_t_3); __pyx_t_3 = 0; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 24, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; } __Pyx_DECREF_SET(__pyx_v_image, __pyx_t_1); __pyx_t_1 = 0;
+025: return self._run(image, run_type=run_type)
__Pyx_XDECREF(__pyx_r); __pyx_t_1 = __Pyx_PyObject_GetAttrStr(__pyx_v_self, __pyx_n_s_run); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 25, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_2 = PyTuple_New(1); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 25, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_INCREF(__pyx_v_image); __Pyx_GIVEREF(__pyx_v_image); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_2, 0, __pyx_v_image)) __PYX_ERR(0, 25, __pyx_L1_error); __pyx_t_3 = __Pyx_PyDict_NewPresized(1); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 25, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); if (PyDict_SetItem(__pyx_t_3, __pyx_n_s_run_type, __pyx_v_run_type) < 0) __PYX_ERR(0, 25, __pyx_L1_error) __pyx_t_5 = __Pyx_PyObject_Call(__pyx_t_1, __pyx_t_2, __pyx_t_3); if (unlikely(!__pyx_t_5)) __PYX_ERR(0, 25, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_5); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __pyx_r = __pyx_t_5; __pyx_t_5 = 0; goto __pyx_L0;
026:
+027: def benchmark(self, image):
/* Python wrapper */
static PyObject *__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_5benchmark(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
); /*proto*/
static PyMethodDef __pyx_mdef_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_5benchmark = {"benchmark", (PyCFunction)(void*)(__Pyx_PyCFunction_FastCallWithKeywords)__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_5benchmark, __Pyx_METH_FASTCALL|METH_KEYWORDS, 0};
static PyObject *__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_5benchmark(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
) {
PyObject *__pyx_v_self = 0;
PyObject *__pyx_v_image = 0;
#if !CYTHON_METH_FASTCALL
CYTHON_UNUSED Py_ssize_t __pyx_nargs;
#endif
CYTHON_UNUSED PyObject *const *__pyx_kwvalues;
PyObject *__pyx_r = 0;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("benchmark (wrapper)", 0);
#if !CYTHON_METH_FASTCALL
#if CYTHON_ASSUME_SAFE_MACROS
__pyx_nargs = PyTuple_GET_SIZE(__pyx_args);
#else
__pyx_nargs = PyTuple_Size(__pyx_args);
if (unlikely((__pyx_nargs < 0))) __PYX_ERR(0, 27, __pyx_L3_error)
#endif
#endif
__pyx_kwvalues = __Pyx_KwValues_FASTCALL(__pyx_args, __pyx_nargs);
{
PyObject **__pyx_pyargnames[] = {&__pyx_n_s_self,&__pyx_n_s_image,0};
PyObject* values[2] = {0,0};
if (__pyx_kwds) {
Py_ssize_t kw_args;
switch (__pyx_nargs) {
case 2: values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
CYTHON_FALLTHROUGH;
case 1: values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
CYTHON_FALLTHROUGH;
case 0: break;
default: goto __pyx_L5_argtuple_error;
}
kw_args = __Pyx_NumKwargs_FASTCALL(__pyx_kwds);
switch (__pyx_nargs) {
case 0:
if (likely((values[0] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_self)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[0]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 27, __pyx_L3_error)
else goto __pyx_L5_argtuple_error;
CYTHON_FALLTHROUGH;
case 1:
if (likely((values[1] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_image)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[1]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 27, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("benchmark", 1, 2, 2, 1); __PYX_ERR(0, 27, __pyx_L3_error)
}
}
if (unlikely(kw_args > 0)) {
const Py_ssize_t kwd_pos_args = __pyx_nargs;
if (unlikely(__Pyx_ParseOptionalKeywords(__pyx_kwds, __pyx_kwvalues, __pyx_pyargnames, 0, values + 0, kwd_pos_args, "benchmark") < 0)) __PYX_ERR(0, 27, __pyx_L3_error)
}
} else if (unlikely(__pyx_nargs != 2)) {
goto __pyx_L5_argtuple_error;
} else {
values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
}
__pyx_v_self = values[0];
__pyx_v_image = values[1];
}
goto __pyx_L4_argument_unpacking_done;
__pyx_L5_argtuple_error:;
__Pyx_RaiseArgtupleInvalid("benchmark", 1, 2, 2, __pyx_nargs); __PYX_ERR(0, 27, __pyx_L3_error)
goto __pyx_L3_error;
__pyx_L3_error:;
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__Pyx_AddTraceback("nanopyx.liquid._le_roberts_cross_gradients.GradientRobertsCross.benchmark", __pyx_clineno, __pyx_lineno, __pyx_filename);
__Pyx_RefNannyFinishContext();
return NULL;
__pyx_L4_argument_unpacking_done:;
__pyx_r = __pyx_pf_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_4benchmark(__pyx_self, __pyx_v_self, __pyx_v_image);
int __pyx_lineno = 0;
const char *__pyx_filename = NULL;
int __pyx_clineno = 0;
/* function exit code */
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
static PyObject *__pyx_pf_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_4benchmark(CYTHON_UNUSED PyObject *__pyx_self, PyObject *__pyx_v_self, PyObject *__pyx_v_image) {
PyObject *__pyx_r = NULL;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("benchmark", 0);
__Pyx_INCREF(__pyx_v_image);
/* … */
/* function exit code */
__pyx_L1_error:;
__Pyx_XDECREF(__pyx_t_1);
__Pyx_XDECREF(__pyx_t_2);
__Pyx_XDECREF(__pyx_t_3);
__Pyx_AddTraceback("nanopyx.liquid._le_roberts_cross_gradients.GradientRobertsCross.benchmark", __pyx_clineno, __pyx_lineno, __pyx_filename);
__pyx_r = NULL;
__pyx_L0:;
__Pyx_XDECREF(__pyx_v_image);
__Pyx_XGIVEREF(__pyx_r);
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
/* … */
__pyx_tuple__28 = PyTuple_Pack(2, __pyx_n_s_self, __pyx_n_s_image); if (unlikely(!__pyx_tuple__28)) __PYX_ERR(0, 27, __pyx_L1_error)
__Pyx_GOTREF(__pyx_tuple__28);
__Pyx_GIVEREF(__pyx_tuple__28);
/* … */
__pyx_t_10 = __Pyx_CyFunction_New(&__pyx_mdef_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_5benchmark, 0, __pyx_n_s_GradientRobertsCross_benchmark, NULL, __pyx_n_s_nanopyx_liquid__le_roberts_cross, __pyx_d, ((PyObject *)__pyx_codeobj__29)); if (unlikely(!__pyx_t_10)) __PYX_ERR(0, 27, __pyx_L1_error)
__Pyx_GOTREF(__pyx_t_10);
__Pyx_INCREF(__pyx_t_10);
PyList_Append(__pyx_t_7, __pyx_t_10);
__Pyx_GIVEREF(__pyx_t_10);
if (__Pyx_SetNameInClass(__pyx_t_9, __pyx_n_s_benchmark, __pyx_t_10) < 0) __PYX_ERR(0, 27, __pyx_L1_error)
__Pyx_DECREF(__pyx_t_10); __pyx_t_10 = 0;
__pyx_codeobj__29 = (PyObject*)__Pyx_PyCode_New(2, 0, 0, 2, 0, CO_OPTIMIZED|CO_NEWLOCALS, __pyx_empty_bytes, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_tuple__28, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_kp_s_src_nanopyx_liquid__le_roberts_c, __pyx_n_s_benchmark, 27, __pyx_empty_bytes); if (unlikely(!__pyx_codeobj__29)) __PYX_ERR(0, 27, __pyx_L1_error)
+028: image = check_image(image)
__Pyx_GetModuleGlobalName(__pyx_t_2, __pyx_n_s_check_image); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_t_3 = NULL; __pyx_t_4 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_2))) { __pyx_t_3 = PyMethod_GET_SELF(__pyx_t_2); if (likely(__pyx_t_3)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_2); __Pyx_INCREF(__pyx_t_3); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_2, function); __pyx_t_4 = 1; } } #endif { PyObject *__pyx_callargs[2] = {__pyx_t_3, __pyx_v_image}; __pyx_t_1 = __Pyx_PyObject_FastCall(__pyx_t_2, __pyx_callargs+1-__pyx_t_4, 1+__pyx_t_4); __Pyx_XDECREF(__pyx_t_3); __pyx_t_3 = 0; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 28, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; } __Pyx_DECREF_SET(__pyx_v_image, __pyx_t_1); __pyx_t_1 = 0;
+029: return super().benchmark(image)
__Pyx_XDECREF(__pyx_r); __pyx_t_2 = __Pyx_CyFunction_GetClassObj(__pyx_self); if (!__pyx_t_2) { PyErr_SetString(PyExc_SystemError, "super(): empty __class__ cell"); __PYX_ERR(0, 29, __pyx_L1_error) } __Pyx_INCREF(__pyx_t_2); __pyx_t_3 = PyTuple_New(2); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 29, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_GIVEREF(__pyx_t_2); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_3, 0, __pyx_t_2)) __PYX_ERR(0, 29, __pyx_L1_error); __Pyx_INCREF(__pyx_v_self); __Pyx_GIVEREF(__pyx_v_self); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_3, 1, __pyx_v_self)) __PYX_ERR(0, 29, __pyx_L1_error); __pyx_t_2 = 0; __pyx_t_2 = __Pyx_PyObject_Call(__pyx_builtin_super, __pyx_t_3, NULL); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 29, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __pyx_t_3 = __Pyx_PyObject_GetAttrStr(__pyx_t_2, __pyx_n_s_benchmark); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 29, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = NULL; __pyx_t_4 = 0; #if CYTHON_UNPACK_METHODS if (likely(PyMethod_Check(__pyx_t_3))) { __pyx_t_2 = PyMethod_GET_SELF(__pyx_t_3); if (likely(__pyx_t_2)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_3); __Pyx_INCREF(__pyx_t_2); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_3, function); __pyx_t_4 = 1; } } #endif { PyObject *__pyx_callargs[2] = {__pyx_t_2, __pyx_v_image}; __pyx_t_1 = __Pyx_PyObject_FastCall(__pyx_t_3, __pyx_callargs+1-__pyx_t_4, 1+__pyx_t_4); __Pyx_XDECREF(__pyx_t_2); __pyx_t_2 = 0; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 29, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; } __pyx_r = __pyx_t_1; __pyx_t_1 = 0; goto __pyx_L0;
030:
031:
032: # tag-start: _le_roberts_cross_gradients.GradientRobertsCross._run_unthreaded
+033: def _run_unthreaded(self, float[:,:,:] image):
/* Python wrapper */
static PyObject *__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_7_run_unthreaded(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
); /*proto*/
static PyMethodDef __pyx_mdef_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_7_run_unthreaded = {"_run_unthreaded", (PyCFunction)(void*)(__Pyx_PyCFunction_FastCallWithKeywords)__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_7_run_unthreaded, __Pyx_METH_FASTCALL|METH_KEYWORDS, 0};
static PyObject *__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_7_run_unthreaded(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
) {
CYTHON_UNUSED PyObject *__pyx_v_self = 0;
__Pyx_memviewslice __pyx_v_image = { 0, 0, { 0 }, { 0 }, { 0 } };
#if !CYTHON_METH_FASTCALL
CYTHON_UNUSED Py_ssize_t __pyx_nargs;
#endif
CYTHON_UNUSED PyObject *const *__pyx_kwvalues;
PyObject *__pyx_r = 0;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("_run_unthreaded (wrapper)", 0);
#if !CYTHON_METH_FASTCALL
#if CYTHON_ASSUME_SAFE_MACROS
__pyx_nargs = PyTuple_GET_SIZE(__pyx_args);
#else
__pyx_nargs = PyTuple_Size(__pyx_args);
if (unlikely((__pyx_nargs < 0))) __PYX_ERR(0, 33, __pyx_L3_error)
#endif
#endif
__pyx_kwvalues = __Pyx_KwValues_FASTCALL(__pyx_args, __pyx_nargs);
{
PyObject **__pyx_pyargnames[] = {&__pyx_n_s_self,&__pyx_n_s_image,0};
PyObject* values[2] = {0,0};
if (__pyx_kwds) {
Py_ssize_t kw_args;
switch (__pyx_nargs) {
case 2: values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
CYTHON_FALLTHROUGH;
case 1: values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
CYTHON_FALLTHROUGH;
case 0: break;
default: goto __pyx_L5_argtuple_error;
}
kw_args = __Pyx_NumKwargs_FASTCALL(__pyx_kwds);
switch (__pyx_nargs) {
case 0:
if (likely((values[0] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_self)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[0]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 33, __pyx_L3_error)
else goto __pyx_L5_argtuple_error;
CYTHON_FALLTHROUGH;
case 1:
if (likely((values[1] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_image)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[1]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 33, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("_run_unthreaded", 1, 2, 2, 1); __PYX_ERR(0, 33, __pyx_L3_error)
}
}
if (unlikely(kw_args > 0)) {
const Py_ssize_t kwd_pos_args = __pyx_nargs;
if (unlikely(__Pyx_ParseOptionalKeywords(__pyx_kwds, __pyx_kwvalues, __pyx_pyargnames, 0, values + 0, kwd_pos_args, "_run_unthreaded") < 0)) __PYX_ERR(0, 33, __pyx_L3_error)
}
} else if (unlikely(__pyx_nargs != 2)) {
goto __pyx_L5_argtuple_error;
} else {
values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
}
__pyx_v_self = values[0];
__pyx_v_image = __Pyx_PyObject_to_MemoryviewSlice_dsdsds_float(values[1], PyBUF_WRITABLE); if (unlikely(!__pyx_v_image.memview)) __PYX_ERR(0, 33, __pyx_L3_error)
}
goto __pyx_L4_argument_unpacking_done;
__pyx_L5_argtuple_error:;
__Pyx_RaiseArgtupleInvalid("_run_unthreaded", 1, 2, 2, __pyx_nargs); __PYX_ERR(0, 33, __pyx_L3_error)
goto __pyx_L3_error;
__pyx_L3_error:;
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__PYX_XCLEAR_MEMVIEW(&__pyx_v_image, 1);
__Pyx_AddTraceback("nanopyx.liquid._le_roberts_cross_gradients.GradientRobertsCross._run_unthreaded", __pyx_clineno, __pyx_lineno, __pyx_filename);
__Pyx_RefNannyFinishContext();
return NULL;
__pyx_L4_argument_unpacking_done:;
__pyx_r = __pyx_pf_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_6_run_unthreaded(__pyx_self, __pyx_v_self, __pyx_v_image);
int __pyx_lineno = 0;
const char *__pyx_filename = NULL;
int __pyx_clineno = 0;
/* function exit code */
__PYX_XCLEAR_MEMVIEW(&__pyx_v_image, 1);
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
static PyObject *__pyx_pf_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_6_run_unthreaded(CYTHON_UNUSED PyObject *__pyx_self, CYTHON_UNUSED PyObject *__pyx_v_self, __Pyx_memviewslice __pyx_v_image) {
int __pyx_v_nFrames;
__Pyx_memviewslice __pyx_v_gradient_col = { 0, 0, { 0 }, { 0 }, { 0 } };
__Pyx_memviewslice __pyx_v_gradient_row = { 0, 0, { 0 }, { 0 }, { 0 } };
int __pyx_v_n;
PyObject *__pyx_r = NULL;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("_run_unthreaded", 0);
/* … */
/* function exit code */
__pyx_L1_error:;
__Pyx_XDECREF(__pyx_t_1);
__Pyx_XDECREF(__pyx_t_2);
__Pyx_XDECREF(__pyx_t_3);
__Pyx_XDECREF(__pyx_t_4);
__PYX_XCLEAR_MEMVIEW(&__pyx_t_6, 1);
__Pyx_AddTraceback("nanopyx.liquid._le_roberts_cross_gradients.GradientRobertsCross._run_unthreaded", __pyx_clineno, __pyx_lineno, __pyx_filename);
__pyx_r = NULL;
__pyx_L0:;
__PYX_XCLEAR_MEMVIEW(&__pyx_v_gradient_col, 1);
__PYX_XCLEAR_MEMVIEW(&__pyx_v_gradient_row, 1);
__Pyx_XGIVEREF(__pyx_r);
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
/* … */
__pyx_tuple__30 = PyTuple_Pack(6, __pyx_n_s_self, __pyx_n_s_image, __pyx_n_s_nFrames, __pyx_n_s_gradient_col, __pyx_n_s_gradient_row, __pyx_n_s_n); if (unlikely(!__pyx_tuple__30)) __PYX_ERR(0, 33, __pyx_L1_error)
__Pyx_GOTREF(__pyx_tuple__30);
__Pyx_GIVEREF(__pyx_tuple__30);
/* … */
__pyx_t_10 = __Pyx_CyFunction_New(&__pyx_mdef_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_7_run_unthreaded, 0, __pyx_n_s_GradientRobertsCross__run_unthre, NULL, __pyx_n_s_nanopyx_liquid__le_roberts_cross, __pyx_d, ((PyObject *)__pyx_codeobj__31)); if (unlikely(!__pyx_t_10)) __PYX_ERR(0, 33, __pyx_L1_error)
__Pyx_GOTREF(__pyx_t_10);
if (__Pyx_SetNameInClass(__pyx_t_9, __pyx_n_s_run_unthreaded, __pyx_t_10) < 0) __PYX_ERR(0, 33, __pyx_L1_error)
__Pyx_DECREF(__pyx_t_10); __pyx_t_10 = 0;
__pyx_codeobj__31 = (PyObject*)__Pyx_PyCode_New(2, 0, 0, 6, 0, CO_OPTIMIZED|CO_NEWLOCALS, __pyx_empty_bytes, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_tuple__30, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_kp_s_src_nanopyx_liquid__le_roberts_c, __pyx_n_s_run_unthreaded, 33, __pyx_empty_bytes); if (unlikely(!__pyx_codeobj__31)) __PYX_ERR(0, 33, __pyx_L1_error)
034:
+035: cdef int nFrames = image.shape[0]
__pyx_v_nFrames = (__pyx_v_image.shape[0]);
+036: cdef float [:,:,:] gradient_col = np.zeros_like(image)
__Pyx_GetModuleGlobalName(__pyx_t_2, __pyx_n_s_np); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 36, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_t_3 = __Pyx_PyObject_GetAttrStr(__pyx_t_2, __pyx_n_s_zeros_like); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 36, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = __pyx_memoryview_fromslice(__pyx_v_image, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 36, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_t_4 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_3))) { __pyx_t_4 = PyMethod_GET_SELF(__pyx_t_3); if (likely(__pyx_t_4)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_3); __Pyx_INCREF(__pyx_t_4); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_3, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[2] = {__pyx_t_4, __pyx_t_2}; __pyx_t_1 = __Pyx_PyObject_FastCall(__pyx_t_3, __pyx_callargs+1-__pyx_t_5, 1+__pyx_t_5); __Pyx_XDECREF(__pyx_t_4); __pyx_t_4 = 0; __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 36, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; } __pyx_t_6 = __Pyx_PyObject_to_MemoryviewSlice_dsdsds_float(__pyx_t_1, PyBUF_WRITABLE); if (unlikely(!__pyx_t_6.memview)) __PYX_ERR(0, 36, __pyx_L1_error) __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_v_gradient_col = __pyx_t_6; __pyx_t_6.memview = NULL; __pyx_t_6.data = NULL;
+037: cdef float [:,:,:] gradient_row = np.zeros_like(image)
__Pyx_GetModuleGlobalName(__pyx_t_3, __pyx_n_s_np); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 37, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_2 = __Pyx_PyObject_GetAttrStr(__pyx_t_3, __pyx_n_s_zeros_like); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 37, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __pyx_t_3 = __pyx_memoryview_fromslice(__pyx_v_image, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 37, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_4 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_2))) { __pyx_t_4 = PyMethod_GET_SELF(__pyx_t_2); if (likely(__pyx_t_4)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_2); __Pyx_INCREF(__pyx_t_4); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_2, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[2] = {__pyx_t_4, __pyx_t_3}; __pyx_t_1 = __Pyx_PyObject_FastCall(__pyx_t_2, __pyx_callargs+1-__pyx_t_5, 1+__pyx_t_5); __Pyx_XDECREF(__pyx_t_4); __pyx_t_4 = 0; __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 37, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; } __pyx_t_6 = __Pyx_PyObject_to_MemoryviewSlice_dsdsds_float(__pyx_t_1, PyBUF_WRITABLE); if (unlikely(!__pyx_t_6.memview)) __PYX_ERR(0, 37, __pyx_L1_error) __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_v_gradient_row = __pyx_t_6; __pyx_t_6.memview = NULL; __pyx_t_6.data = NULL;
038:
039: cdef int n
+040: with nogil:
{
#ifdef WITH_THREAD
PyThreadState *_save;
_save = NULL;
Py_UNBLOCK_THREADS
__Pyx_FastGIL_Remember();
#endif
/*try:*/ {
/* … */
/*finally:*/ {
/*normal exit:*/{
#ifdef WITH_THREAD
__Pyx_FastGIL_Forget();
Py_BLOCK_THREADS
#endif
goto __pyx_L5;
}
__pyx_L5:;
}
}
+041: for n in range(nFrames):
__pyx_t_5 = __pyx_v_nFrames;
__pyx_t_7 = __pyx_t_5;
for (__pyx_t_8 = 0; __pyx_t_8 < __pyx_t_7; __pyx_t_8+=1) {
__pyx_v_n = __pyx_t_8;
+042: _c_gradient_roberts_cross(&image[n,0,0], &gradient_col[n,0,0], &gradient_row[n,0,0], image.shape[1], image.shape[2])
__pyx_t_9 = __pyx_v_n;
__pyx_t_10 = 0;
__pyx_t_11 = 0;
__pyx_t_12 = __pyx_v_n;
__pyx_t_13 = 0;
__pyx_t_14 = 0;
__pyx_t_15 = __pyx_v_n;
__pyx_t_16 = 0;
__pyx_t_17 = 0;
_c_gradient_roberts_cross((&(*((float *) ( /* dim=2 */ (( /* dim=1 */ (( /* dim=0 */ (__pyx_v_image.data + __pyx_t_9 * __pyx_v_image.strides[0]) ) + __pyx_t_10 * __pyx_v_image.strides[1]) ) + __pyx_t_11 * __pyx_v_image.strides[2]) )))), (&(*((float *) ( /* dim=2 */ (( /* dim=1 */ (( /* dim=0 */ (__pyx_v_gradient_col.data + __pyx_t_12 * __pyx_v_gradient_col.strides[0]) ) + __pyx_t_13 * __pyx_v_gradient_col.strides[1]) ) + __pyx_t_14 * __pyx_v_gradient_col.strides[2]) )))), (&(*((float *) ( /* dim=2 */ (( /* dim=1 */ (( /* dim=0 */ (__pyx_v_gradient_row.data + __pyx_t_15 * __pyx_v_gradient_row.strides[0]) ) + __pyx_t_16 * __pyx_v_gradient_row.strides[1]) ) + __pyx_t_17 * __pyx_v_gradient_row.strides[2]) )))), (__pyx_v_image.shape[1]), (__pyx_v_image.shape[2]));
}
}
043:
+044: return gradient_col, gradient_row
__Pyx_XDECREF(__pyx_r); __pyx_t_1 = __pyx_memoryview_fromslice(__pyx_v_gradient_col, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 44, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_2 = __pyx_memoryview_fromslice(__pyx_v_gradient_row, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 44, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_t_3 = PyTuple_New(2); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 44, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_GIVEREF(__pyx_t_1); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_3, 0, __pyx_t_1)) __PYX_ERR(0, 44, __pyx_L1_error); __Pyx_GIVEREF(__pyx_t_2); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_3, 1, __pyx_t_2)) __PYX_ERR(0, 44, __pyx_L1_error); __pyx_t_1 = 0; __pyx_t_2 = 0; __pyx_r = __pyx_t_3; __pyx_t_3 = 0; goto __pyx_L0;
045: # tag-end
046:
047:
048: # tag-copy: _le_roberts_cross_gradients.GradientRobertsCross._run_unthreaded; replace("_run_unthreaded", "_run_threaded"); replace("range(nFrames)", "prange(nFrames)")
+049: def _run_threaded(self, float[:,:,:] image):
/* Python wrapper */
static PyObject *__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_9_run_threaded(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
); /*proto*/
static PyMethodDef __pyx_mdef_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_9_run_threaded = {"_run_threaded", (PyCFunction)(void*)(__Pyx_PyCFunction_FastCallWithKeywords)__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_9_run_threaded, __Pyx_METH_FASTCALL|METH_KEYWORDS, 0};
static PyObject *__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_9_run_threaded(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
) {
CYTHON_UNUSED PyObject *__pyx_v_self = 0;
__Pyx_memviewslice __pyx_v_image = { 0, 0, { 0 }, { 0 }, { 0 } };
#if !CYTHON_METH_FASTCALL
CYTHON_UNUSED Py_ssize_t __pyx_nargs;
#endif
CYTHON_UNUSED PyObject *const *__pyx_kwvalues;
PyObject *__pyx_r = 0;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("_run_threaded (wrapper)", 0);
#if !CYTHON_METH_FASTCALL
#if CYTHON_ASSUME_SAFE_MACROS
__pyx_nargs = PyTuple_GET_SIZE(__pyx_args);
#else
__pyx_nargs = PyTuple_Size(__pyx_args);
if (unlikely((__pyx_nargs < 0))) __PYX_ERR(0, 49, __pyx_L3_error)
#endif
#endif
__pyx_kwvalues = __Pyx_KwValues_FASTCALL(__pyx_args, __pyx_nargs);
{
PyObject **__pyx_pyargnames[] = {&__pyx_n_s_self,&__pyx_n_s_image,0};
PyObject* values[2] = {0,0};
if (__pyx_kwds) {
Py_ssize_t kw_args;
switch (__pyx_nargs) {
case 2: values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
CYTHON_FALLTHROUGH;
case 1: values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
CYTHON_FALLTHROUGH;
case 0: break;
default: goto __pyx_L5_argtuple_error;
}
kw_args = __Pyx_NumKwargs_FASTCALL(__pyx_kwds);
switch (__pyx_nargs) {
case 0:
if (likely((values[0] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_self)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[0]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 49, __pyx_L3_error)
else goto __pyx_L5_argtuple_error;
CYTHON_FALLTHROUGH;
case 1:
if (likely((values[1] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_image)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[1]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 49, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("_run_threaded", 1, 2, 2, 1); __PYX_ERR(0, 49, __pyx_L3_error)
}
}
if (unlikely(kw_args > 0)) {
const Py_ssize_t kwd_pos_args = __pyx_nargs;
if (unlikely(__Pyx_ParseOptionalKeywords(__pyx_kwds, __pyx_kwvalues, __pyx_pyargnames, 0, values + 0, kwd_pos_args, "_run_threaded") < 0)) __PYX_ERR(0, 49, __pyx_L3_error)
}
} else if (unlikely(__pyx_nargs != 2)) {
goto __pyx_L5_argtuple_error;
} else {
values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
}
__pyx_v_self = values[0];
__pyx_v_image = __Pyx_PyObject_to_MemoryviewSlice_dsdsds_float(values[1], PyBUF_WRITABLE); if (unlikely(!__pyx_v_image.memview)) __PYX_ERR(0, 49, __pyx_L3_error)
}
goto __pyx_L4_argument_unpacking_done;
__pyx_L5_argtuple_error:;
__Pyx_RaiseArgtupleInvalid("_run_threaded", 1, 2, 2, __pyx_nargs); __PYX_ERR(0, 49, __pyx_L3_error)
goto __pyx_L3_error;
__pyx_L3_error:;
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__PYX_XCLEAR_MEMVIEW(&__pyx_v_image, 1);
__Pyx_AddTraceback("nanopyx.liquid._le_roberts_cross_gradients.GradientRobertsCross._run_threaded", __pyx_clineno, __pyx_lineno, __pyx_filename);
__Pyx_RefNannyFinishContext();
return NULL;
__pyx_L4_argument_unpacking_done:;
__pyx_r = __pyx_pf_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_8_run_threaded(__pyx_self, __pyx_v_self, __pyx_v_image);
int __pyx_lineno = 0;
const char *__pyx_filename = NULL;
int __pyx_clineno = 0;
/* function exit code */
__PYX_XCLEAR_MEMVIEW(&__pyx_v_image, 1);
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
static PyObject *__pyx_pf_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_8_run_threaded(CYTHON_UNUSED PyObject *__pyx_self, CYTHON_UNUSED PyObject *__pyx_v_self, __Pyx_memviewslice __pyx_v_image) {
CYTHON_UNUSED int __pyx_v_nFrames;
__Pyx_memviewslice __pyx_v_gradient_col = { 0, 0, { 0 }, { 0 }, { 0 } };
__Pyx_memviewslice __pyx_v_gradient_row = { 0, 0, { 0 }, { 0 }, { 0 } };
int __pyx_v_n;
PyObject *__pyx_r = NULL;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("_run_threaded", 0);
/* … */
/* function exit code */
__pyx_L1_error:;
__Pyx_XDECREF(__pyx_t_1);
__Pyx_XDECREF(__pyx_t_2);
__Pyx_XDECREF(__pyx_t_3);
__Pyx_XDECREF(__pyx_t_4);
__PYX_XCLEAR_MEMVIEW(&__pyx_t_6, 1);
__Pyx_AddTraceback("nanopyx.liquid._le_roberts_cross_gradients.GradientRobertsCross._run_threaded", __pyx_clineno, __pyx_lineno, __pyx_filename);
__pyx_r = NULL;
__pyx_L0:;
__PYX_XCLEAR_MEMVIEW(&__pyx_v_gradient_col, 1);
__PYX_XCLEAR_MEMVIEW(&__pyx_v_gradient_row, 1);
__Pyx_XGIVEREF(__pyx_r);
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
/* … */
__pyx_t_10 = __Pyx_CyFunction_New(&__pyx_mdef_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_9_run_threaded, 0, __pyx_n_s_GradientRobertsCross__run_thread, NULL, __pyx_n_s_nanopyx_liquid__le_roberts_cross, __pyx_d, ((PyObject *)__pyx_codeobj__32)); if (unlikely(!__pyx_t_10)) __PYX_ERR(0, 49, __pyx_L1_error)
__Pyx_GOTREF(__pyx_t_10);
if (__Pyx_SetNameInClass(__pyx_t_9, __pyx_n_s_run_threaded, __pyx_t_10) < 0) __PYX_ERR(0, 49, __pyx_L1_error)
__Pyx_DECREF(__pyx_t_10); __pyx_t_10 = 0;
/* … */
__pyx_codeobj__32 = (PyObject*)__Pyx_PyCode_New(2, 0, 0, 6, 0, CO_OPTIMIZED|CO_NEWLOCALS, __pyx_empty_bytes, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_tuple__30, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_kp_s_src_nanopyx_liquid__le_roberts_c, __pyx_n_s_run_threaded, 49, __pyx_empty_bytes); if (unlikely(!__pyx_codeobj__32)) __PYX_ERR(0, 49, __pyx_L1_error)
050:
+051: cdef int nFrames = image.shape[0]
__pyx_v_nFrames = (__pyx_v_image.shape[0]);
+052: cdef float [:,:,:] gradient_col = np.zeros_like(image)
__Pyx_GetModuleGlobalName(__pyx_t_2, __pyx_n_s_np); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 52, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_t_3 = __Pyx_PyObject_GetAttrStr(__pyx_t_2, __pyx_n_s_zeros_like); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 52, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = __pyx_memoryview_fromslice(__pyx_v_image, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 52, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_t_4 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_3))) { __pyx_t_4 = PyMethod_GET_SELF(__pyx_t_3); if (likely(__pyx_t_4)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_3); __Pyx_INCREF(__pyx_t_4); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_3, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[2] = {__pyx_t_4, __pyx_t_2}; __pyx_t_1 = __Pyx_PyObject_FastCall(__pyx_t_3, __pyx_callargs+1-__pyx_t_5, 1+__pyx_t_5); __Pyx_XDECREF(__pyx_t_4); __pyx_t_4 = 0; __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 52, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; } __pyx_t_6 = __Pyx_PyObject_to_MemoryviewSlice_dsdsds_float(__pyx_t_1, PyBUF_WRITABLE); if (unlikely(!__pyx_t_6.memview)) __PYX_ERR(0, 52, __pyx_L1_error) __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_v_gradient_col = __pyx_t_6; __pyx_t_6.memview = NULL; __pyx_t_6.data = NULL;
+053: cdef float [:,:,:] gradient_row = np.zeros_like(image)
__Pyx_GetModuleGlobalName(__pyx_t_3, __pyx_n_s_np); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 53, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_2 = __Pyx_PyObject_GetAttrStr(__pyx_t_3, __pyx_n_s_zeros_like); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 53, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __pyx_t_3 = __pyx_memoryview_fromslice(__pyx_v_image, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 53, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_4 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_2))) { __pyx_t_4 = PyMethod_GET_SELF(__pyx_t_2); if (likely(__pyx_t_4)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_2); __Pyx_INCREF(__pyx_t_4); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_2, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[2] = {__pyx_t_4, __pyx_t_3}; __pyx_t_1 = __Pyx_PyObject_FastCall(__pyx_t_2, __pyx_callargs+1-__pyx_t_5, 1+__pyx_t_5); __Pyx_XDECREF(__pyx_t_4); __pyx_t_4 = 0; __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 53, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; } __pyx_t_6 = __Pyx_PyObject_to_MemoryviewSlice_dsdsds_float(__pyx_t_1, PyBUF_WRITABLE); if (unlikely(!__pyx_t_6.memview)) __PYX_ERR(0, 53, __pyx_L1_error) __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_v_gradient_row = __pyx_t_6; __pyx_t_6.memview = NULL; __pyx_t_6.data = NULL;
054:
055: cdef int n
+056: with nogil:
{
#ifdef WITH_THREAD
PyThreadState *_save;
_save = NULL;
Py_UNBLOCK_THREADS
__Pyx_FastGIL_Remember();
#endif
/*try:*/ {
/* … */
/*finally:*/ {
/*normal exit:*/{
#ifdef WITH_THREAD
__Pyx_FastGIL_Forget();
Py_BLOCK_THREADS
#endif
goto __pyx_L5;
}
__pyx_L5:;
}
}
+057: for n in prange(nFrames):
__pyx_t_5 = __pyx_v_nFrames;
{
#if ((defined(__APPLE__) || defined(__OSX__)) && (defined(__GNUC__) && (__GNUC__ > 2 || (__GNUC__ == 2 && (__GNUC_MINOR__ > 95)))))
#undef likely
#undef unlikely
#define likely(x) (x)
#define unlikely(x) (x)
#endif
__pyx_t_8 = (__pyx_t_5 - 0 + 1 - 1/abs(1)) / 1;
if (__pyx_t_8 > 0)
{
#ifdef _OPENMP
#pragma omp parallel
#endif /* _OPENMP */
{
#ifdef _OPENMP
#pragma omp for firstprivate(__pyx_v_n) lastprivate(__pyx_v_n)
#endif /* _OPENMP */
for (__pyx_t_7 = 0; __pyx_t_7 < __pyx_t_8; __pyx_t_7++){
{
__pyx_v_n = (int)(0 + 1 * __pyx_t_7);
+058: _c_gradient_roberts_cross(&image[n,0,0], &gradient_col[n,0,0], &gradient_row[n,0,0], image.shape[1], image.shape[2])
__pyx_t_9 = __pyx_v_n;
__pyx_t_10 = 0;
__pyx_t_11 = 0;
__pyx_t_12 = __pyx_v_n;
__pyx_t_13 = 0;
__pyx_t_14 = 0;
__pyx_t_15 = __pyx_v_n;
__pyx_t_16 = 0;
__pyx_t_17 = 0;
_c_gradient_roberts_cross((&(*((float *) ( /* dim=2 */ (( /* dim=1 */ (( /* dim=0 */ (__pyx_v_image.data + __pyx_t_9 * __pyx_v_image.strides[0]) ) + __pyx_t_10 * __pyx_v_image.strides[1]) ) + __pyx_t_11 * __pyx_v_image.strides[2]) )))), (&(*((float *) ( /* dim=2 */ (( /* dim=1 */ (( /* dim=0 */ (__pyx_v_gradient_col.data + __pyx_t_12 * __pyx_v_gradient_col.strides[0]) ) + __pyx_t_13 * __pyx_v_gradient_col.strides[1]) ) + __pyx_t_14 * __pyx_v_gradient_col.strides[2]) )))), (&(*((float *) ( /* dim=2 */ (( /* dim=1 */ (( /* dim=0 */ (__pyx_v_gradient_row.data + __pyx_t_15 * __pyx_v_gradient_row.strides[0]) ) + __pyx_t_16 * __pyx_v_gradient_row.strides[1]) ) + __pyx_t_17 * __pyx_v_gradient_row.strides[2]) )))), (__pyx_v_image.shape[1]), (__pyx_v_image.shape[2]));
}
}
}
}
}
#if ((defined(__APPLE__) || defined(__OSX__)) && (defined(__GNUC__) && (__GNUC__ > 2 || (__GNUC__ == 2 && (__GNUC_MINOR__ > 95)))))
#undef likely
#undef unlikely
#define likely(x) __builtin_expect(!!(x), 1)
#define unlikely(x) __builtin_expect(!!(x), 0)
#endif
}
059:
+060: return gradient_col, gradient_row
__Pyx_XDECREF(__pyx_r); __pyx_t_1 = __pyx_memoryview_fromslice(__pyx_v_gradient_col, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 60, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_2 = __pyx_memoryview_fromslice(__pyx_v_gradient_row, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 60, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_t_3 = PyTuple_New(2); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 60, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_GIVEREF(__pyx_t_1); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_3, 0, __pyx_t_1)) __PYX_ERR(0, 60, __pyx_L1_error); __Pyx_GIVEREF(__pyx_t_2); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_3, 1, __pyx_t_2)) __PYX_ERR(0, 60, __pyx_L1_error); __pyx_t_1 = 0; __pyx_t_2 = 0; __pyx_r = __pyx_t_3; __pyx_t_3 = 0; goto __pyx_L0;
061: # tag-end
062:
063:
064: # tag-copy: _le_roberts_cross_gradients.GradientRobertsCross._run_unthreaded; replace("_run_unthreaded", "_run_threaded_static"); replace("range(nFrames)", 'prange(nFrames, schedule="static")')
+065: def _run_threaded_static(self, float[:,:,:] image):
/* Python wrapper */
static PyObject *__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_11_run_threaded_static(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
); /*proto*/
static PyMethodDef __pyx_mdef_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_11_run_threaded_static = {"_run_threaded_static", (PyCFunction)(void*)(__Pyx_PyCFunction_FastCallWithKeywords)__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_11_run_threaded_static, __Pyx_METH_FASTCALL|METH_KEYWORDS, 0};
static PyObject *__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_11_run_threaded_static(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
) {
CYTHON_UNUSED PyObject *__pyx_v_self = 0;
__Pyx_memviewslice __pyx_v_image = { 0, 0, { 0 }, { 0 }, { 0 } };
#if !CYTHON_METH_FASTCALL
CYTHON_UNUSED Py_ssize_t __pyx_nargs;
#endif
CYTHON_UNUSED PyObject *const *__pyx_kwvalues;
PyObject *__pyx_r = 0;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("_run_threaded_static (wrapper)", 0);
#if !CYTHON_METH_FASTCALL
#if CYTHON_ASSUME_SAFE_MACROS
__pyx_nargs = PyTuple_GET_SIZE(__pyx_args);
#else
__pyx_nargs = PyTuple_Size(__pyx_args);
if (unlikely((__pyx_nargs < 0))) __PYX_ERR(0, 65, __pyx_L3_error)
#endif
#endif
__pyx_kwvalues = __Pyx_KwValues_FASTCALL(__pyx_args, __pyx_nargs);
{
PyObject **__pyx_pyargnames[] = {&__pyx_n_s_self,&__pyx_n_s_image,0};
PyObject* values[2] = {0,0};
if (__pyx_kwds) {
Py_ssize_t kw_args;
switch (__pyx_nargs) {
case 2: values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
CYTHON_FALLTHROUGH;
case 1: values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
CYTHON_FALLTHROUGH;
case 0: break;
default: goto __pyx_L5_argtuple_error;
}
kw_args = __Pyx_NumKwargs_FASTCALL(__pyx_kwds);
switch (__pyx_nargs) {
case 0:
if (likely((values[0] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_self)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[0]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 65, __pyx_L3_error)
else goto __pyx_L5_argtuple_error;
CYTHON_FALLTHROUGH;
case 1:
if (likely((values[1] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_image)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[1]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 65, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("_run_threaded_static", 1, 2, 2, 1); __PYX_ERR(0, 65, __pyx_L3_error)
}
}
if (unlikely(kw_args > 0)) {
const Py_ssize_t kwd_pos_args = __pyx_nargs;
if (unlikely(__Pyx_ParseOptionalKeywords(__pyx_kwds, __pyx_kwvalues, __pyx_pyargnames, 0, values + 0, kwd_pos_args, "_run_threaded_static") < 0)) __PYX_ERR(0, 65, __pyx_L3_error)
}
} else if (unlikely(__pyx_nargs != 2)) {
goto __pyx_L5_argtuple_error;
} else {
values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
}
__pyx_v_self = values[0];
__pyx_v_image = __Pyx_PyObject_to_MemoryviewSlice_dsdsds_float(values[1], PyBUF_WRITABLE); if (unlikely(!__pyx_v_image.memview)) __PYX_ERR(0, 65, __pyx_L3_error)
}
goto __pyx_L4_argument_unpacking_done;
__pyx_L5_argtuple_error:;
__Pyx_RaiseArgtupleInvalid("_run_threaded_static", 1, 2, 2, __pyx_nargs); __PYX_ERR(0, 65, __pyx_L3_error)
goto __pyx_L3_error;
__pyx_L3_error:;
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__PYX_XCLEAR_MEMVIEW(&__pyx_v_image, 1);
__Pyx_AddTraceback("nanopyx.liquid._le_roberts_cross_gradients.GradientRobertsCross._run_threaded_static", __pyx_clineno, __pyx_lineno, __pyx_filename);
__Pyx_RefNannyFinishContext();
return NULL;
__pyx_L4_argument_unpacking_done:;
__pyx_r = __pyx_pf_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_10_run_threaded_static(__pyx_self, __pyx_v_self, __pyx_v_image);
int __pyx_lineno = 0;
const char *__pyx_filename = NULL;
int __pyx_clineno = 0;
/* function exit code */
__PYX_XCLEAR_MEMVIEW(&__pyx_v_image, 1);
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
static PyObject *__pyx_pf_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_10_run_threaded_static(CYTHON_UNUSED PyObject *__pyx_self, CYTHON_UNUSED PyObject *__pyx_v_self, __Pyx_memviewslice __pyx_v_image) {
CYTHON_UNUSED int __pyx_v_nFrames;
__Pyx_memviewslice __pyx_v_gradient_col = { 0, 0, { 0 }, { 0 }, { 0 } };
__Pyx_memviewslice __pyx_v_gradient_row = { 0, 0, { 0 }, { 0 }, { 0 } };
int __pyx_v_n;
PyObject *__pyx_r = NULL;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("_run_threaded_static", 0);
/* … */
/* function exit code */
__pyx_L1_error:;
__Pyx_XDECREF(__pyx_t_1);
__Pyx_XDECREF(__pyx_t_2);
__Pyx_XDECREF(__pyx_t_3);
__Pyx_XDECREF(__pyx_t_4);
__PYX_XCLEAR_MEMVIEW(&__pyx_t_6, 1);
__Pyx_AddTraceback("nanopyx.liquid._le_roberts_cross_gradients.GradientRobertsCross._run_threaded_static", __pyx_clineno, __pyx_lineno, __pyx_filename);
__pyx_r = NULL;
__pyx_L0:;
__PYX_XCLEAR_MEMVIEW(&__pyx_v_gradient_col, 1);
__PYX_XCLEAR_MEMVIEW(&__pyx_v_gradient_row, 1);
__Pyx_XGIVEREF(__pyx_r);
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
/* … */
__pyx_t_10 = __Pyx_CyFunction_New(&__pyx_mdef_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_11_run_threaded_static, 0, __pyx_n_s_GradientRobertsCross__run_thread_2, NULL, __pyx_n_s_nanopyx_liquid__le_roberts_cross, __pyx_d, ((PyObject *)__pyx_codeobj__33)); if (unlikely(!__pyx_t_10)) __PYX_ERR(0, 65, __pyx_L1_error)
__Pyx_GOTREF(__pyx_t_10);
if (__Pyx_SetNameInClass(__pyx_t_9, __pyx_n_s_run_threaded_static, __pyx_t_10) < 0) __PYX_ERR(0, 65, __pyx_L1_error)
__Pyx_DECREF(__pyx_t_10); __pyx_t_10 = 0;
/* … */
__pyx_codeobj__33 = (PyObject*)__Pyx_PyCode_New(2, 0, 0, 6, 0, CO_OPTIMIZED|CO_NEWLOCALS, __pyx_empty_bytes, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_tuple__30, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_kp_s_src_nanopyx_liquid__le_roberts_c, __pyx_n_s_run_threaded_static, 65, __pyx_empty_bytes); if (unlikely(!__pyx_codeobj__33)) __PYX_ERR(0, 65, __pyx_L1_error)
066:
+067: cdef int nFrames = image.shape[0]
__pyx_v_nFrames = (__pyx_v_image.shape[0]);
+068: cdef float [:,:,:] gradient_col = np.zeros_like(image)
__Pyx_GetModuleGlobalName(__pyx_t_2, __pyx_n_s_np); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 68, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_t_3 = __Pyx_PyObject_GetAttrStr(__pyx_t_2, __pyx_n_s_zeros_like); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 68, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = __pyx_memoryview_fromslice(__pyx_v_image, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 68, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_t_4 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_3))) { __pyx_t_4 = PyMethod_GET_SELF(__pyx_t_3); if (likely(__pyx_t_4)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_3); __Pyx_INCREF(__pyx_t_4); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_3, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[2] = {__pyx_t_4, __pyx_t_2}; __pyx_t_1 = __Pyx_PyObject_FastCall(__pyx_t_3, __pyx_callargs+1-__pyx_t_5, 1+__pyx_t_5); __Pyx_XDECREF(__pyx_t_4); __pyx_t_4 = 0; __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 68, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; } __pyx_t_6 = __Pyx_PyObject_to_MemoryviewSlice_dsdsds_float(__pyx_t_1, PyBUF_WRITABLE); if (unlikely(!__pyx_t_6.memview)) __PYX_ERR(0, 68, __pyx_L1_error) __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_v_gradient_col = __pyx_t_6; __pyx_t_6.memview = NULL; __pyx_t_6.data = NULL;
+069: cdef float [:,:,:] gradient_row = np.zeros_like(image)
__Pyx_GetModuleGlobalName(__pyx_t_3, __pyx_n_s_np); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 69, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_2 = __Pyx_PyObject_GetAttrStr(__pyx_t_3, __pyx_n_s_zeros_like); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 69, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __pyx_t_3 = __pyx_memoryview_fromslice(__pyx_v_image, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 69, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_4 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_2))) { __pyx_t_4 = PyMethod_GET_SELF(__pyx_t_2); if (likely(__pyx_t_4)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_2); __Pyx_INCREF(__pyx_t_4); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_2, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[2] = {__pyx_t_4, __pyx_t_3}; __pyx_t_1 = __Pyx_PyObject_FastCall(__pyx_t_2, __pyx_callargs+1-__pyx_t_5, 1+__pyx_t_5); __Pyx_XDECREF(__pyx_t_4); __pyx_t_4 = 0; __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 69, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; } __pyx_t_6 = __Pyx_PyObject_to_MemoryviewSlice_dsdsds_float(__pyx_t_1, PyBUF_WRITABLE); if (unlikely(!__pyx_t_6.memview)) __PYX_ERR(0, 69, __pyx_L1_error) __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_v_gradient_row = __pyx_t_6; __pyx_t_6.memview = NULL; __pyx_t_6.data = NULL;
070:
071: cdef int n
+072: with nogil:
{
#ifdef WITH_THREAD
PyThreadState *_save;
_save = NULL;
Py_UNBLOCK_THREADS
__Pyx_FastGIL_Remember();
#endif
/*try:*/ {
/* … */
/*finally:*/ {
/*normal exit:*/{
#ifdef WITH_THREAD
__Pyx_FastGIL_Forget();
Py_BLOCK_THREADS
#endif
goto __pyx_L5;
}
__pyx_L5:;
}
}
+073: for n in prange(nFrames, schedule="static"):
__pyx_t_5 = __pyx_v_nFrames;
{
#if ((defined(__APPLE__) || defined(__OSX__)) && (defined(__GNUC__) && (__GNUC__ > 2 || (__GNUC__ == 2 && (__GNUC_MINOR__ > 95)))))
#undef likely
#undef unlikely
#define likely(x) (x)
#define unlikely(x) (x)
#endif
__pyx_t_8 = (__pyx_t_5 - 0 + 1 - 1/abs(1)) / 1;
if (__pyx_t_8 > 0)
{
#ifdef _OPENMP
#pragma omp parallel
#endif /* _OPENMP */
{
#ifdef _OPENMP
#pragma omp for firstprivate(__pyx_v_n) lastprivate(__pyx_v_n) schedule(static)
#endif /* _OPENMP */
for (__pyx_t_7 = 0; __pyx_t_7 < __pyx_t_8; __pyx_t_7++){
{
__pyx_v_n = (int)(0 + 1 * __pyx_t_7);
+074: _c_gradient_roberts_cross(&image[n,0,0], &gradient_col[n,0,0], &gradient_row[n,0,0], image.shape[1], image.shape[2])
__pyx_t_9 = __pyx_v_n;
__pyx_t_10 = 0;
__pyx_t_11 = 0;
__pyx_t_12 = __pyx_v_n;
__pyx_t_13 = 0;
__pyx_t_14 = 0;
__pyx_t_15 = __pyx_v_n;
__pyx_t_16 = 0;
__pyx_t_17 = 0;
_c_gradient_roberts_cross((&(*((float *) ( /* dim=2 */ (( /* dim=1 */ (( /* dim=0 */ (__pyx_v_image.data + __pyx_t_9 * __pyx_v_image.strides[0]) ) + __pyx_t_10 * __pyx_v_image.strides[1]) ) + __pyx_t_11 * __pyx_v_image.strides[2]) )))), (&(*((float *) ( /* dim=2 */ (( /* dim=1 */ (( /* dim=0 */ (__pyx_v_gradient_col.data + __pyx_t_12 * __pyx_v_gradient_col.strides[0]) ) + __pyx_t_13 * __pyx_v_gradient_col.strides[1]) ) + __pyx_t_14 * __pyx_v_gradient_col.strides[2]) )))), (&(*((float *) ( /* dim=2 */ (( /* dim=1 */ (( /* dim=0 */ (__pyx_v_gradient_row.data + __pyx_t_15 * __pyx_v_gradient_row.strides[0]) ) + __pyx_t_16 * __pyx_v_gradient_row.strides[1]) ) + __pyx_t_17 * __pyx_v_gradient_row.strides[2]) )))), (__pyx_v_image.shape[1]), (__pyx_v_image.shape[2]));
}
}
}
}
}
#if ((defined(__APPLE__) || defined(__OSX__)) && (defined(__GNUC__) && (__GNUC__ > 2 || (__GNUC__ == 2 && (__GNUC_MINOR__ > 95)))))
#undef likely
#undef unlikely
#define likely(x) __builtin_expect(!!(x), 1)
#define unlikely(x) __builtin_expect(!!(x), 0)
#endif
}
075:
+076: return gradient_col, gradient_row
__Pyx_XDECREF(__pyx_r); __pyx_t_1 = __pyx_memoryview_fromslice(__pyx_v_gradient_col, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 76, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_2 = __pyx_memoryview_fromslice(__pyx_v_gradient_row, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 76, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_t_3 = PyTuple_New(2); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 76, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_GIVEREF(__pyx_t_1); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_3, 0, __pyx_t_1)) __PYX_ERR(0, 76, __pyx_L1_error); __Pyx_GIVEREF(__pyx_t_2); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_3, 1, __pyx_t_2)) __PYX_ERR(0, 76, __pyx_L1_error); __pyx_t_1 = 0; __pyx_t_2 = 0; __pyx_r = __pyx_t_3; __pyx_t_3 = 0; goto __pyx_L0;
077: # tag-end
078:
079:
080: # tag-copy: _le_roberts_cross_gradients.GradientRobertsCross._run_unthreaded; replace("_run_unthreaded", "_run_threaded_dynamic"); replace("range(nFrames)", 'prange(nFrames, schedule="dynamic")')
+081: def _run_threaded_dynamic(self, float[:,:,:] image):
/* Python wrapper */
static PyObject *__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_13_run_threaded_dynamic(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
); /*proto*/
static PyMethodDef __pyx_mdef_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_13_run_threaded_dynamic = {"_run_threaded_dynamic", (PyCFunction)(void*)(__Pyx_PyCFunction_FastCallWithKeywords)__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_13_run_threaded_dynamic, __Pyx_METH_FASTCALL|METH_KEYWORDS, 0};
static PyObject *__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_13_run_threaded_dynamic(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
) {
CYTHON_UNUSED PyObject *__pyx_v_self = 0;
__Pyx_memviewslice __pyx_v_image = { 0, 0, { 0 }, { 0 }, { 0 } };
#if !CYTHON_METH_FASTCALL
CYTHON_UNUSED Py_ssize_t __pyx_nargs;
#endif
CYTHON_UNUSED PyObject *const *__pyx_kwvalues;
PyObject *__pyx_r = 0;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("_run_threaded_dynamic (wrapper)", 0);
#if !CYTHON_METH_FASTCALL
#if CYTHON_ASSUME_SAFE_MACROS
__pyx_nargs = PyTuple_GET_SIZE(__pyx_args);
#else
__pyx_nargs = PyTuple_Size(__pyx_args);
if (unlikely((__pyx_nargs < 0))) __PYX_ERR(0, 81, __pyx_L3_error)
#endif
#endif
__pyx_kwvalues = __Pyx_KwValues_FASTCALL(__pyx_args, __pyx_nargs);
{
PyObject **__pyx_pyargnames[] = {&__pyx_n_s_self,&__pyx_n_s_image,0};
PyObject* values[2] = {0,0};
if (__pyx_kwds) {
Py_ssize_t kw_args;
switch (__pyx_nargs) {
case 2: values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
CYTHON_FALLTHROUGH;
case 1: values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
CYTHON_FALLTHROUGH;
case 0: break;
default: goto __pyx_L5_argtuple_error;
}
kw_args = __Pyx_NumKwargs_FASTCALL(__pyx_kwds);
switch (__pyx_nargs) {
case 0:
if (likely((values[0] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_self)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[0]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 81, __pyx_L3_error)
else goto __pyx_L5_argtuple_error;
CYTHON_FALLTHROUGH;
case 1:
if (likely((values[1] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_image)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[1]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 81, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("_run_threaded_dynamic", 1, 2, 2, 1); __PYX_ERR(0, 81, __pyx_L3_error)
}
}
if (unlikely(kw_args > 0)) {
const Py_ssize_t kwd_pos_args = __pyx_nargs;
if (unlikely(__Pyx_ParseOptionalKeywords(__pyx_kwds, __pyx_kwvalues, __pyx_pyargnames, 0, values + 0, kwd_pos_args, "_run_threaded_dynamic") < 0)) __PYX_ERR(0, 81, __pyx_L3_error)
}
} else if (unlikely(__pyx_nargs != 2)) {
goto __pyx_L5_argtuple_error;
} else {
values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
}
__pyx_v_self = values[0];
__pyx_v_image = __Pyx_PyObject_to_MemoryviewSlice_dsdsds_float(values[1], PyBUF_WRITABLE); if (unlikely(!__pyx_v_image.memview)) __PYX_ERR(0, 81, __pyx_L3_error)
}
goto __pyx_L4_argument_unpacking_done;
__pyx_L5_argtuple_error:;
__Pyx_RaiseArgtupleInvalid("_run_threaded_dynamic", 1, 2, 2, __pyx_nargs); __PYX_ERR(0, 81, __pyx_L3_error)
goto __pyx_L3_error;
__pyx_L3_error:;
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__PYX_XCLEAR_MEMVIEW(&__pyx_v_image, 1);
__Pyx_AddTraceback("nanopyx.liquid._le_roberts_cross_gradients.GradientRobertsCross._run_threaded_dynamic", __pyx_clineno, __pyx_lineno, __pyx_filename);
__Pyx_RefNannyFinishContext();
return NULL;
__pyx_L4_argument_unpacking_done:;
__pyx_r = __pyx_pf_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_12_run_threaded_dynamic(__pyx_self, __pyx_v_self, __pyx_v_image);
int __pyx_lineno = 0;
const char *__pyx_filename = NULL;
int __pyx_clineno = 0;
/* function exit code */
__PYX_XCLEAR_MEMVIEW(&__pyx_v_image, 1);
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
static PyObject *__pyx_pf_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_12_run_threaded_dynamic(CYTHON_UNUSED PyObject *__pyx_self, CYTHON_UNUSED PyObject *__pyx_v_self, __Pyx_memviewslice __pyx_v_image) {
CYTHON_UNUSED int __pyx_v_nFrames;
__Pyx_memviewslice __pyx_v_gradient_col = { 0, 0, { 0 }, { 0 }, { 0 } };
__Pyx_memviewslice __pyx_v_gradient_row = { 0, 0, { 0 }, { 0 }, { 0 } };
int __pyx_v_n;
PyObject *__pyx_r = NULL;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("_run_threaded_dynamic", 0);
/* … */
/* function exit code */
__pyx_L1_error:;
__Pyx_XDECREF(__pyx_t_1);
__Pyx_XDECREF(__pyx_t_2);
__Pyx_XDECREF(__pyx_t_3);
__Pyx_XDECREF(__pyx_t_4);
__PYX_XCLEAR_MEMVIEW(&__pyx_t_6, 1);
__Pyx_AddTraceback("nanopyx.liquid._le_roberts_cross_gradients.GradientRobertsCross._run_threaded_dynamic", __pyx_clineno, __pyx_lineno, __pyx_filename);
__pyx_r = NULL;
__pyx_L0:;
__PYX_XCLEAR_MEMVIEW(&__pyx_v_gradient_col, 1);
__PYX_XCLEAR_MEMVIEW(&__pyx_v_gradient_row, 1);
__Pyx_XGIVEREF(__pyx_r);
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
/* … */
__pyx_t_10 = __Pyx_CyFunction_New(&__pyx_mdef_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_13_run_threaded_dynamic, 0, __pyx_n_s_GradientRobertsCross__run_thread_3, NULL, __pyx_n_s_nanopyx_liquid__le_roberts_cross, __pyx_d, ((PyObject *)__pyx_codeobj__34)); if (unlikely(!__pyx_t_10)) __PYX_ERR(0, 81, __pyx_L1_error)
__Pyx_GOTREF(__pyx_t_10);
if (__Pyx_SetNameInClass(__pyx_t_9, __pyx_n_s_run_threaded_dynamic, __pyx_t_10) < 0) __PYX_ERR(0, 81, __pyx_L1_error)
__Pyx_DECREF(__pyx_t_10); __pyx_t_10 = 0;
/* … */
__pyx_codeobj__34 = (PyObject*)__Pyx_PyCode_New(2, 0, 0, 6, 0, CO_OPTIMIZED|CO_NEWLOCALS, __pyx_empty_bytes, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_tuple__30, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_kp_s_src_nanopyx_liquid__le_roberts_c, __pyx_n_s_run_threaded_dynamic, 81, __pyx_empty_bytes); if (unlikely(!__pyx_codeobj__34)) __PYX_ERR(0, 81, __pyx_L1_error)
082:
+083: cdef int nFrames = image.shape[0]
__pyx_v_nFrames = (__pyx_v_image.shape[0]);
+084: cdef float [:,:,:] gradient_col = np.zeros_like(image)
__Pyx_GetModuleGlobalName(__pyx_t_2, __pyx_n_s_np); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 84, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_t_3 = __Pyx_PyObject_GetAttrStr(__pyx_t_2, __pyx_n_s_zeros_like); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 84, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = __pyx_memoryview_fromslice(__pyx_v_image, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 84, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_t_4 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_3))) { __pyx_t_4 = PyMethod_GET_SELF(__pyx_t_3); if (likely(__pyx_t_4)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_3); __Pyx_INCREF(__pyx_t_4); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_3, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[2] = {__pyx_t_4, __pyx_t_2}; __pyx_t_1 = __Pyx_PyObject_FastCall(__pyx_t_3, __pyx_callargs+1-__pyx_t_5, 1+__pyx_t_5); __Pyx_XDECREF(__pyx_t_4); __pyx_t_4 = 0; __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 84, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; } __pyx_t_6 = __Pyx_PyObject_to_MemoryviewSlice_dsdsds_float(__pyx_t_1, PyBUF_WRITABLE); if (unlikely(!__pyx_t_6.memview)) __PYX_ERR(0, 84, __pyx_L1_error) __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_v_gradient_col = __pyx_t_6; __pyx_t_6.memview = NULL; __pyx_t_6.data = NULL;
+085: cdef float [:,:,:] gradient_row = np.zeros_like(image)
__Pyx_GetModuleGlobalName(__pyx_t_3, __pyx_n_s_np); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 85, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_2 = __Pyx_PyObject_GetAttrStr(__pyx_t_3, __pyx_n_s_zeros_like); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 85, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __pyx_t_3 = __pyx_memoryview_fromslice(__pyx_v_image, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 85, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_4 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_2))) { __pyx_t_4 = PyMethod_GET_SELF(__pyx_t_2); if (likely(__pyx_t_4)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_2); __Pyx_INCREF(__pyx_t_4); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_2, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[2] = {__pyx_t_4, __pyx_t_3}; __pyx_t_1 = __Pyx_PyObject_FastCall(__pyx_t_2, __pyx_callargs+1-__pyx_t_5, 1+__pyx_t_5); __Pyx_XDECREF(__pyx_t_4); __pyx_t_4 = 0; __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 85, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; } __pyx_t_6 = __Pyx_PyObject_to_MemoryviewSlice_dsdsds_float(__pyx_t_1, PyBUF_WRITABLE); if (unlikely(!__pyx_t_6.memview)) __PYX_ERR(0, 85, __pyx_L1_error) __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_v_gradient_row = __pyx_t_6; __pyx_t_6.memview = NULL; __pyx_t_6.data = NULL;
086:
087: cdef int n
+088: with nogil:
{
#ifdef WITH_THREAD
PyThreadState *_save;
_save = NULL;
Py_UNBLOCK_THREADS
__Pyx_FastGIL_Remember();
#endif
/*try:*/ {
/* … */
/*finally:*/ {
/*normal exit:*/{
#ifdef WITH_THREAD
__Pyx_FastGIL_Forget();
Py_BLOCK_THREADS
#endif
goto __pyx_L5;
}
__pyx_L5:;
}
}
+089: for n in prange(nFrames, schedule="dynamic"):
__pyx_t_5 = __pyx_v_nFrames;
{
#if ((defined(__APPLE__) || defined(__OSX__)) && (defined(__GNUC__) && (__GNUC__ > 2 || (__GNUC__ == 2 && (__GNUC_MINOR__ > 95)))))
#undef likely
#undef unlikely
#define likely(x) (x)
#define unlikely(x) (x)
#endif
__pyx_t_8 = (__pyx_t_5 - 0 + 1 - 1/abs(1)) / 1;
if (__pyx_t_8 > 0)
{
#ifdef _OPENMP
#pragma omp parallel
#endif /* _OPENMP */
{
#ifdef _OPENMP
#pragma omp for firstprivate(__pyx_v_n) lastprivate(__pyx_v_n) schedule(dynamic)
#endif /* _OPENMP */
for (__pyx_t_7 = 0; __pyx_t_7 < __pyx_t_8; __pyx_t_7++){
{
__pyx_v_n = (int)(0 + 1 * __pyx_t_7);
+090: _c_gradient_roberts_cross(&image[n,0,0], &gradient_col[n,0,0], &gradient_row[n,0,0], image.shape[1], image.shape[2])
__pyx_t_9 = __pyx_v_n;
__pyx_t_10 = 0;
__pyx_t_11 = 0;
__pyx_t_12 = __pyx_v_n;
__pyx_t_13 = 0;
__pyx_t_14 = 0;
__pyx_t_15 = __pyx_v_n;
__pyx_t_16 = 0;
__pyx_t_17 = 0;
_c_gradient_roberts_cross((&(*((float *) ( /* dim=2 */ (( /* dim=1 */ (( /* dim=0 */ (__pyx_v_image.data + __pyx_t_9 * __pyx_v_image.strides[0]) ) + __pyx_t_10 * __pyx_v_image.strides[1]) ) + __pyx_t_11 * __pyx_v_image.strides[2]) )))), (&(*((float *) ( /* dim=2 */ (( /* dim=1 */ (( /* dim=0 */ (__pyx_v_gradient_col.data + __pyx_t_12 * __pyx_v_gradient_col.strides[0]) ) + __pyx_t_13 * __pyx_v_gradient_col.strides[1]) ) + __pyx_t_14 * __pyx_v_gradient_col.strides[2]) )))), (&(*((float *) ( /* dim=2 */ (( /* dim=1 */ (( /* dim=0 */ (__pyx_v_gradient_row.data + __pyx_t_15 * __pyx_v_gradient_row.strides[0]) ) + __pyx_t_16 * __pyx_v_gradient_row.strides[1]) ) + __pyx_t_17 * __pyx_v_gradient_row.strides[2]) )))), (__pyx_v_image.shape[1]), (__pyx_v_image.shape[2]));
}
}
}
}
}
#if ((defined(__APPLE__) || defined(__OSX__)) && (defined(__GNUC__) && (__GNUC__ > 2 || (__GNUC__ == 2 && (__GNUC_MINOR__ > 95)))))
#undef likely
#undef unlikely
#define likely(x) __builtin_expect(!!(x), 1)
#define unlikely(x) __builtin_expect(!!(x), 0)
#endif
}
091:
+092: return gradient_col, gradient_row
__Pyx_XDECREF(__pyx_r); __pyx_t_1 = __pyx_memoryview_fromslice(__pyx_v_gradient_col, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 92, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_2 = __pyx_memoryview_fromslice(__pyx_v_gradient_row, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 92, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_t_3 = PyTuple_New(2); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 92, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_GIVEREF(__pyx_t_1); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_3, 0, __pyx_t_1)) __PYX_ERR(0, 92, __pyx_L1_error); __Pyx_GIVEREF(__pyx_t_2); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_3, 1, __pyx_t_2)) __PYX_ERR(0, 92, __pyx_L1_error); __pyx_t_1 = 0; __pyx_t_2 = 0; __pyx_r = __pyx_t_3; __pyx_t_3 = 0; goto __pyx_L0;
093: # tag-end
094:
095:
096: # tag-copy: _le_roberts_cross_gradients.GradientRobertsCross._run_unthreaded; replace("_run_unthreaded", "_run_threaded_guided"); replace("range(nFrames)", 'prange(nFrames, schedule="guided")')
+097: def _run_threaded_guided(self, float[:,:,:] image):
/* Python wrapper */
static PyObject *__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_15_run_threaded_guided(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
); /*proto*/
static PyMethodDef __pyx_mdef_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_15_run_threaded_guided = {"_run_threaded_guided", (PyCFunction)(void*)(__Pyx_PyCFunction_FastCallWithKeywords)__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_15_run_threaded_guided, __Pyx_METH_FASTCALL|METH_KEYWORDS, 0};
static PyObject *__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_15_run_threaded_guided(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
) {
CYTHON_UNUSED PyObject *__pyx_v_self = 0;
__Pyx_memviewslice __pyx_v_image = { 0, 0, { 0 }, { 0 }, { 0 } };
#if !CYTHON_METH_FASTCALL
CYTHON_UNUSED Py_ssize_t __pyx_nargs;
#endif
CYTHON_UNUSED PyObject *const *__pyx_kwvalues;
PyObject *__pyx_r = 0;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("_run_threaded_guided (wrapper)", 0);
#if !CYTHON_METH_FASTCALL
#if CYTHON_ASSUME_SAFE_MACROS
__pyx_nargs = PyTuple_GET_SIZE(__pyx_args);
#else
__pyx_nargs = PyTuple_Size(__pyx_args);
if (unlikely((__pyx_nargs < 0))) __PYX_ERR(0, 97, __pyx_L3_error)
#endif
#endif
__pyx_kwvalues = __Pyx_KwValues_FASTCALL(__pyx_args, __pyx_nargs);
{
PyObject **__pyx_pyargnames[] = {&__pyx_n_s_self,&__pyx_n_s_image,0};
PyObject* values[2] = {0,0};
if (__pyx_kwds) {
Py_ssize_t kw_args;
switch (__pyx_nargs) {
case 2: values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
CYTHON_FALLTHROUGH;
case 1: values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
CYTHON_FALLTHROUGH;
case 0: break;
default: goto __pyx_L5_argtuple_error;
}
kw_args = __Pyx_NumKwargs_FASTCALL(__pyx_kwds);
switch (__pyx_nargs) {
case 0:
if (likely((values[0] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_self)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[0]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 97, __pyx_L3_error)
else goto __pyx_L5_argtuple_error;
CYTHON_FALLTHROUGH;
case 1:
if (likely((values[1] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_image)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[1]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 97, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("_run_threaded_guided", 1, 2, 2, 1); __PYX_ERR(0, 97, __pyx_L3_error)
}
}
if (unlikely(kw_args > 0)) {
const Py_ssize_t kwd_pos_args = __pyx_nargs;
if (unlikely(__Pyx_ParseOptionalKeywords(__pyx_kwds, __pyx_kwvalues, __pyx_pyargnames, 0, values + 0, kwd_pos_args, "_run_threaded_guided") < 0)) __PYX_ERR(0, 97, __pyx_L3_error)
}
} else if (unlikely(__pyx_nargs != 2)) {
goto __pyx_L5_argtuple_error;
} else {
values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
}
__pyx_v_self = values[0];
__pyx_v_image = __Pyx_PyObject_to_MemoryviewSlice_dsdsds_float(values[1], PyBUF_WRITABLE); if (unlikely(!__pyx_v_image.memview)) __PYX_ERR(0, 97, __pyx_L3_error)
}
goto __pyx_L4_argument_unpacking_done;
__pyx_L5_argtuple_error:;
__Pyx_RaiseArgtupleInvalid("_run_threaded_guided", 1, 2, 2, __pyx_nargs); __PYX_ERR(0, 97, __pyx_L3_error)
goto __pyx_L3_error;
__pyx_L3_error:;
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__PYX_XCLEAR_MEMVIEW(&__pyx_v_image, 1);
__Pyx_AddTraceback("nanopyx.liquid._le_roberts_cross_gradients.GradientRobertsCross._run_threaded_guided", __pyx_clineno, __pyx_lineno, __pyx_filename);
__Pyx_RefNannyFinishContext();
return NULL;
__pyx_L4_argument_unpacking_done:;
__pyx_r = __pyx_pf_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_14_run_threaded_guided(__pyx_self, __pyx_v_self, __pyx_v_image);
int __pyx_lineno = 0;
const char *__pyx_filename = NULL;
int __pyx_clineno = 0;
/* function exit code */
__PYX_XCLEAR_MEMVIEW(&__pyx_v_image, 1);
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
static PyObject *__pyx_pf_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_14_run_threaded_guided(CYTHON_UNUSED PyObject *__pyx_self, CYTHON_UNUSED PyObject *__pyx_v_self, __Pyx_memviewslice __pyx_v_image) {
CYTHON_UNUSED int __pyx_v_nFrames;
__Pyx_memviewslice __pyx_v_gradient_col = { 0, 0, { 0 }, { 0 }, { 0 } };
__Pyx_memviewslice __pyx_v_gradient_row = { 0, 0, { 0 }, { 0 }, { 0 } };
int __pyx_v_n;
PyObject *__pyx_r = NULL;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("_run_threaded_guided", 0);
/* … */
/* function exit code */
__pyx_L1_error:;
__Pyx_XDECREF(__pyx_t_1);
__Pyx_XDECREF(__pyx_t_2);
__Pyx_XDECREF(__pyx_t_3);
__Pyx_XDECREF(__pyx_t_4);
__PYX_XCLEAR_MEMVIEW(&__pyx_t_6, 1);
__Pyx_AddTraceback("nanopyx.liquid._le_roberts_cross_gradients.GradientRobertsCross._run_threaded_guided", __pyx_clineno, __pyx_lineno, __pyx_filename);
__pyx_r = NULL;
__pyx_L0:;
__PYX_XCLEAR_MEMVIEW(&__pyx_v_gradient_col, 1);
__PYX_XCLEAR_MEMVIEW(&__pyx_v_gradient_row, 1);
__Pyx_XGIVEREF(__pyx_r);
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
/* … */
__pyx_t_10 = __Pyx_CyFunction_New(&__pyx_mdef_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_15_run_threaded_guided, 0, __pyx_n_s_GradientRobertsCross__run_thread_4, NULL, __pyx_n_s_nanopyx_liquid__le_roberts_cross, __pyx_d, ((PyObject *)__pyx_codeobj__35)); if (unlikely(!__pyx_t_10)) __PYX_ERR(0, 97, __pyx_L1_error)
__Pyx_GOTREF(__pyx_t_10);
if (__Pyx_SetNameInClass(__pyx_t_9, __pyx_n_s_run_threaded_guided, __pyx_t_10) < 0) __PYX_ERR(0, 97, __pyx_L1_error)
__Pyx_DECREF(__pyx_t_10); __pyx_t_10 = 0;
/* … */
__pyx_codeobj__35 = (PyObject*)__Pyx_PyCode_New(2, 0, 0, 6, 0, CO_OPTIMIZED|CO_NEWLOCALS, __pyx_empty_bytes, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_tuple__30, __pyx_empty_tuple, __pyx_empty_tuple, __pyx_kp_s_src_nanopyx_liquid__le_roberts_c, __pyx_n_s_run_threaded_guided, 97, __pyx_empty_bytes); if (unlikely(!__pyx_codeobj__35)) __PYX_ERR(0, 97, __pyx_L1_error)
098:
+099: cdef int nFrames = image.shape[0]
__pyx_v_nFrames = (__pyx_v_image.shape[0]);
+100: cdef float [:,:,:] gradient_col = np.zeros_like(image)
__Pyx_GetModuleGlobalName(__pyx_t_2, __pyx_n_s_np); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 100, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_t_3 = __Pyx_PyObject_GetAttrStr(__pyx_t_2, __pyx_n_s_zeros_like); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 100, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = __pyx_memoryview_fromslice(__pyx_v_image, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 100, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_t_4 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_3))) { __pyx_t_4 = PyMethod_GET_SELF(__pyx_t_3); if (likely(__pyx_t_4)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_3); __Pyx_INCREF(__pyx_t_4); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_3, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[2] = {__pyx_t_4, __pyx_t_2}; __pyx_t_1 = __Pyx_PyObject_FastCall(__pyx_t_3, __pyx_callargs+1-__pyx_t_5, 1+__pyx_t_5); __Pyx_XDECREF(__pyx_t_4); __pyx_t_4 = 0; __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 100, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; } __pyx_t_6 = __Pyx_PyObject_to_MemoryviewSlice_dsdsds_float(__pyx_t_1, PyBUF_WRITABLE); if (unlikely(!__pyx_t_6.memview)) __PYX_ERR(0, 100, __pyx_L1_error) __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_v_gradient_col = __pyx_t_6; __pyx_t_6.memview = NULL; __pyx_t_6.data = NULL;
+101: cdef float [:,:,:] gradient_row = np.zeros_like(image)
__Pyx_GetModuleGlobalName(__pyx_t_3, __pyx_n_s_np); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 101, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_2 = __Pyx_PyObject_GetAttrStr(__pyx_t_3, __pyx_n_s_zeros_like); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 101, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __pyx_t_3 = __pyx_memoryview_fromslice(__pyx_v_image, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 101, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_4 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_2))) { __pyx_t_4 = PyMethod_GET_SELF(__pyx_t_2); if (likely(__pyx_t_4)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_2); __Pyx_INCREF(__pyx_t_4); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_2, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[2] = {__pyx_t_4, __pyx_t_3}; __pyx_t_1 = __Pyx_PyObject_FastCall(__pyx_t_2, __pyx_callargs+1-__pyx_t_5, 1+__pyx_t_5); __Pyx_XDECREF(__pyx_t_4); __pyx_t_4 = 0; __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 101, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; } __pyx_t_6 = __Pyx_PyObject_to_MemoryviewSlice_dsdsds_float(__pyx_t_1, PyBUF_WRITABLE); if (unlikely(!__pyx_t_6.memview)) __PYX_ERR(0, 101, __pyx_L1_error) __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_v_gradient_row = __pyx_t_6; __pyx_t_6.memview = NULL; __pyx_t_6.data = NULL;
102:
103: cdef int n
+104: with nogil:
{
#ifdef WITH_THREAD
PyThreadState *_save;
_save = NULL;
Py_UNBLOCK_THREADS
__Pyx_FastGIL_Remember();
#endif
/*try:*/ {
/* … */
/*finally:*/ {
/*normal exit:*/{
#ifdef WITH_THREAD
__Pyx_FastGIL_Forget();
Py_BLOCK_THREADS
#endif
goto __pyx_L5;
}
__pyx_L5:;
}
}
+105: for n in prange(nFrames, schedule="guided"):
__pyx_t_5 = __pyx_v_nFrames;
{
#if ((defined(__APPLE__) || defined(__OSX__)) && (defined(__GNUC__) && (__GNUC__ > 2 || (__GNUC__ == 2 && (__GNUC_MINOR__ > 95)))))
#undef likely
#undef unlikely
#define likely(x) (x)
#define unlikely(x) (x)
#endif
__pyx_t_8 = (__pyx_t_5 - 0 + 1 - 1/abs(1)) / 1;
if (__pyx_t_8 > 0)
{
#ifdef _OPENMP
#pragma omp parallel
#endif /* _OPENMP */
{
#ifdef _OPENMP
#pragma omp for firstprivate(__pyx_v_n) lastprivate(__pyx_v_n) schedule(guided)
#endif /* _OPENMP */
for (__pyx_t_7 = 0; __pyx_t_7 < __pyx_t_8; __pyx_t_7++){
{
__pyx_v_n = (int)(0 + 1 * __pyx_t_7);
+106: _c_gradient_roberts_cross(&image[n,0,0], &gradient_col[n,0,0], &gradient_row[n,0,0], image.shape[1], image.shape[2])
__pyx_t_9 = __pyx_v_n;
__pyx_t_10 = 0;
__pyx_t_11 = 0;
__pyx_t_12 = __pyx_v_n;
__pyx_t_13 = 0;
__pyx_t_14 = 0;
__pyx_t_15 = __pyx_v_n;
__pyx_t_16 = 0;
__pyx_t_17 = 0;
_c_gradient_roberts_cross((&(*((float *) ( /* dim=2 */ (( /* dim=1 */ (( /* dim=0 */ (__pyx_v_image.data + __pyx_t_9 * __pyx_v_image.strides[0]) ) + __pyx_t_10 * __pyx_v_image.strides[1]) ) + __pyx_t_11 * __pyx_v_image.strides[2]) )))), (&(*((float *) ( /* dim=2 */ (( /* dim=1 */ (( /* dim=0 */ (__pyx_v_gradient_col.data + __pyx_t_12 * __pyx_v_gradient_col.strides[0]) ) + __pyx_t_13 * __pyx_v_gradient_col.strides[1]) ) + __pyx_t_14 * __pyx_v_gradient_col.strides[2]) )))), (&(*((float *) ( /* dim=2 */ (( /* dim=1 */ (( /* dim=0 */ (__pyx_v_gradient_row.data + __pyx_t_15 * __pyx_v_gradient_row.strides[0]) ) + __pyx_t_16 * __pyx_v_gradient_row.strides[1]) ) + __pyx_t_17 * __pyx_v_gradient_row.strides[2]) )))), (__pyx_v_image.shape[1]), (__pyx_v_image.shape[2]));
}
}
}
}
}
#if ((defined(__APPLE__) || defined(__OSX__)) && (defined(__GNUC__) && (__GNUC__ > 2 || (__GNUC__ == 2 && (__GNUC_MINOR__ > 95)))))
#undef likely
#undef unlikely
#define likely(x) __builtin_expect(!!(x), 1)
#define unlikely(x) __builtin_expect(!!(x), 0)
#endif
}
107:
+108: return gradient_col, gradient_row
__Pyx_XDECREF(__pyx_r); __pyx_t_1 = __pyx_memoryview_fromslice(__pyx_v_gradient_col, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 108, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_2 = __pyx_memoryview_fromslice(__pyx_v_gradient_row, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 108, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_t_3 = PyTuple_New(2); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 108, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_GIVEREF(__pyx_t_1); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_3, 0, __pyx_t_1)) __PYX_ERR(0, 108, __pyx_L1_error); __Pyx_GIVEREF(__pyx_t_2); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_3, 1, __pyx_t_2)) __PYX_ERR(0, 108, __pyx_L1_error); __pyx_t_1 = 0; __pyx_t_2 = 0; __pyx_r = __pyx_t_3; __pyx_t_3 = 0; goto __pyx_L0;
109: # tag-end
110:
+111: def _run_opencl(self, float[:,:,:] image, dict device):
/* Python wrapper */
static PyObject *__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_17_run_opencl(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
); /*proto*/
static PyMethodDef __pyx_mdef_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_17_run_opencl = {"_run_opencl", (PyCFunction)(void*)(__Pyx_PyCFunction_FastCallWithKeywords)__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_17_run_opencl, __Pyx_METH_FASTCALL|METH_KEYWORDS, 0};
static PyObject *__pyx_pw_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_17_run_opencl(PyObject *__pyx_self,
#if CYTHON_METH_FASTCALL
PyObject *const *__pyx_args, Py_ssize_t __pyx_nargs, PyObject *__pyx_kwds
#else
PyObject *__pyx_args, PyObject *__pyx_kwds
#endif
) {
PyObject *__pyx_v_self = 0;
__Pyx_memviewslice __pyx_v_image = { 0, 0, { 0 }, { 0 }, { 0 } };
PyObject *__pyx_v_device = 0;
#if !CYTHON_METH_FASTCALL
CYTHON_UNUSED Py_ssize_t __pyx_nargs;
#endif
CYTHON_UNUSED PyObject *const *__pyx_kwvalues;
PyObject *__pyx_r = 0;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("_run_opencl (wrapper)", 0);
#if !CYTHON_METH_FASTCALL
#if CYTHON_ASSUME_SAFE_MACROS
__pyx_nargs = PyTuple_GET_SIZE(__pyx_args);
#else
__pyx_nargs = PyTuple_Size(__pyx_args);
if (unlikely((__pyx_nargs < 0))) __PYX_ERR(0, 111, __pyx_L3_error)
#endif
#endif
__pyx_kwvalues = __Pyx_KwValues_FASTCALL(__pyx_args, __pyx_nargs);
{
PyObject **__pyx_pyargnames[] = {&__pyx_n_s_self,&__pyx_n_s_image,&__pyx_n_s_device,0};
PyObject* values[3] = {0,0,0};
if (__pyx_kwds) {
Py_ssize_t kw_args;
switch (__pyx_nargs) {
case 3: values[2] = __Pyx_Arg_FASTCALL(__pyx_args, 2);
CYTHON_FALLTHROUGH;
case 2: values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
CYTHON_FALLTHROUGH;
case 1: values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
CYTHON_FALLTHROUGH;
case 0: break;
default: goto __pyx_L5_argtuple_error;
}
kw_args = __Pyx_NumKwargs_FASTCALL(__pyx_kwds);
switch (__pyx_nargs) {
case 0:
if (likely((values[0] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_self)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[0]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 111, __pyx_L3_error)
else goto __pyx_L5_argtuple_error;
CYTHON_FALLTHROUGH;
case 1:
if (likely((values[1] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_image)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[1]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 111, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("_run_opencl", 1, 3, 3, 1); __PYX_ERR(0, 111, __pyx_L3_error)
}
CYTHON_FALLTHROUGH;
case 2:
if (likely((values[2] = __Pyx_GetKwValue_FASTCALL(__pyx_kwds, __pyx_kwvalues, __pyx_n_s_device)) != 0)) {
(void)__Pyx_Arg_NewRef_FASTCALL(values[2]);
kw_args--;
}
else if (unlikely(PyErr_Occurred())) __PYX_ERR(0, 111, __pyx_L3_error)
else {
__Pyx_RaiseArgtupleInvalid("_run_opencl", 1, 3, 3, 2); __PYX_ERR(0, 111, __pyx_L3_error)
}
}
if (unlikely(kw_args > 0)) {
const Py_ssize_t kwd_pos_args = __pyx_nargs;
if (unlikely(__Pyx_ParseOptionalKeywords(__pyx_kwds, __pyx_kwvalues, __pyx_pyargnames, 0, values + 0, kwd_pos_args, "_run_opencl") < 0)) __PYX_ERR(0, 111, __pyx_L3_error)
}
} else if (unlikely(__pyx_nargs != 3)) {
goto __pyx_L5_argtuple_error;
} else {
values[0] = __Pyx_Arg_FASTCALL(__pyx_args, 0);
values[1] = __Pyx_Arg_FASTCALL(__pyx_args, 1);
values[2] = __Pyx_Arg_FASTCALL(__pyx_args, 2);
}
__pyx_v_self = values[0];
__pyx_v_image = __Pyx_PyObject_to_MemoryviewSlice_dsdsds_float(values[1], PyBUF_WRITABLE); if (unlikely(!__pyx_v_image.memview)) __PYX_ERR(0, 111, __pyx_L3_error)
__pyx_v_device = ((PyObject*)values[2]);
}
goto __pyx_L4_argument_unpacking_done;
__pyx_L5_argtuple_error:;
__Pyx_RaiseArgtupleInvalid("_run_opencl", 1, 3, 3, __pyx_nargs); __PYX_ERR(0, 111, __pyx_L3_error)
goto __pyx_L3_error;
__pyx_L3_error:;
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__PYX_XCLEAR_MEMVIEW(&__pyx_v_image, 1);
__Pyx_AddTraceback("nanopyx.liquid._le_roberts_cross_gradients.GradientRobertsCross._run_opencl", __pyx_clineno, __pyx_lineno, __pyx_filename);
__Pyx_RefNannyFinishContext();
return NULL;
__pyx_L4_argument_unpacking_done:;
if (unlikely(!__Pyx_ArgTypeTest(((PyObject *)__pyx_v_device), (&PyDict_Type), 1, "device", 1))) __PYX_ERR(0, 111, __pyx_L1_error)
__pyx_r = __pyx_pf_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_16_run_opencl(__pyx_self, __pyx_v_self, __pyx_v_image, __pyx_v_device);
int __pyx_lineno = 0;
const char *__pyx_filename = NULL;
int __pyx_clineno = 0;
/* function exit code */
goto __pyx_L0;
__pyx_L1_error:;
__pyx_r = NULL;
__pyx_L0:;
__PYX_XCLEAR_MEMVIEW(&__pyx_v_image, 1);
{
Py_ssize_t __pyx_temp;
for (__pyx_temp=0; __pyx_temp < (Py_ssize_t)(sizeof(values)/sizeof(values[0])); ++__pyx_temp) {
__Pyx_Arg_XDECREF_FASTCALL(values[__pyx_temp]);
}
}
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
static PyObject *__pyx_pf_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_16_run_opencl(CYTHON_UNUSED PyObject *__pyx_self, PyObject *__pyx_v_self, __Pyx_memviewslice __pyx_v_image, PyObject *__pyx_v_device) {
PyObject *__pyx_v_cl_ctx = NULL;
PyObject *__pyx_v_dc = NULL;
PyObject *__pyx_v_cl_queue = NULL;
CYTHON_UNUSED int __pyx_v_nFrames;
int __pyx_v_row;
int __pyx_v_col;
__Pyx_memviewslice __pyx_v_gradient_col = { 0, 0, { 0 }, { 0 }, { 0 } };
__Pyx_memviewslice __pyx_v_gradient_row = { 0, 0, { 0 }, { 0 }, { 0 } };
PyObject *__pyx_v_max_slices = NULL;
PyObject *__pyx_v_mf = NULL;
PyObject *__pyx_v_input_opencl = NULL;
PyObject *__pyx_v_output_opencl_col = NULL;
PyObject *__pyx_v_output_opencl_row = NULL;
PyObject *__pyx_v_code = NULL;
PyObject *__pyx_v_prg = NULL;
PyObject *__pyx_v_knl = NULL;
PyObject *__pyx_v_i = NULL;
PyObject *__pyx_v_n_slices = NULL;
PyObject *__pyx_r = NULL;
__Pyx_RefNannyDeclarations
__Pyx_RefNannySetupContext("_run_opencl", 0);
/* … */
/* function exit code */
__pyx_L1_error:;
__Pyx_XDECREF(__pyx_t_1);
__Pyx_XDECREF(__pyx_t_2);
__Pyx_XDECREF(__pyx_t_3);
__Pyx_XDECREF(__pyx_t_4);
__PYX_XCLEAR_MEMVIEW(&__pyx_t_6, 1);
__PYX_XCLEAR_MEMVIEW(&__pyx_t_7, 1);
__Pyx_XDECREF(__pyx_t_9);
__Pyx_XDECREF(__pyx_t_12);
__Pyx_XDECREF(__pyx_t_13);
__Pyx_XDECREF(__pyx_t_14);
__Pyx_XDECREF(__pyx_t_15);
__Pyx_AddTraceback("nanopyx.liquid._le_roberts_cross_gradients.GradientRobertsCross._run_opencl", __pyx_clineno, __pyx_lineno, __pyx_filename);
__pyx_r = NULL;
__pyx_L0:;
__Pyx_XDECREF(__pyx_v_cl_ctx);
__Pyx_XDECREF(__pyx_v_dc);
__Pyx_XDECREF(__pyx_v_cl_queue);
__PYX_XCLEAR_MEMVIEW(&__pyx_v_gradient_col, 1);
__PYX_XCLEAR_MEMVIEW(&__pyx_v_gradient_row, 1);
__Pyx_XDECREF(__pyx_v_max_slices);
__Pyx_XDECREF(__pyx_v_mf);
__Pyx_XDECREF(__pyx_v_input_opencl);
__Pyx_XDECREF(__pyx_v_output_opencl_col);
__Pyx_XDECREF(__pyx_v_output_opencl_row);
__Pyx_XDECREF(__pyx_v_code);
__Pyx_XDECREF(__pyx_v_prg);
__Pyx_XDECREF(__pyx_v_knl);
__Pyx_XDECREF(__pyx_v_i);
__Pyx_XDECREF(__pyx_v_n_slices);
__Pyx_XGIVEREF(__pyx_r);
__Pyx_RefNannyFinishContext();
return __pyx_r;
}
/* … */
__pyx_tuple__36 = PyTuple_Pack(21, __pyx_n_s_self, __pyx_n_s_image, __pyx_n_s_device, __pyx_n_s_cl_ctx, __pyx_n_s_dc, __pyx_n_s_cl_queue, __pyx_n_s_nFrames, __pyx_n_s_row, __pyx_n_s_col, __pyx_n_s_gradient_col, __pyx_n_s_gradient_row, __pyx_n_s_max_slices, __pyx_n_s_mf, __pyx_n_s_input_opencl, __pyx_n_s_output_opencl_col, __pyx_n_s_output_opencl_row, __pyx_n_s_code, __pyx_n_s_prg, __pyx_n_s_knl, __pyx_n_s_i, __pyx_n_s_n_slices); if (unlikely(!__pyx_tuple__36)) __PYX_ERR(0, 111, __pyx_L1_error)
__Pyx_GOTREF(__pyx_tuple__36);
__Pyx_GIVEREF(__pyx_tuple__36);
/* … */
__pyx_t_10 = __Pyx_CyFunction_New(&__pyx_mdef_7nanopyx_6liquid_27_le_roberts_cross_gradients_20GradientRobertsCross_17_run_opencl, 0, __pyx_n_s_GradientRobertsCross__run_opencl, NULL, __pyx_n_s_nanopyx_liquid__le_roberts_cross, __pyx_d, ((PyObject *)__pyx_codeobj__37)); if (unlikely(!__pyx_t_10)) __PYX_ERR(0, 111, __pyx_L1_error)
__Pyx_GOTREF(__pyx_t_10);
if (__Pyx_SetNameInClass(__pyx_t_9, __pyx_n_s_run_opencl, __pyx_t_10) < 0) __PYX_ERR(0, 111, __pyx_L1_error)
__Pyx_DECREF(__pyx_t_10); __pyx_t_10 = 0;
112:
113: # QUEUE AND CONTEXT
+114: cl_ctx = cl.Context([device['device']])
__Pyx_GetModuleGlobalName(__pyx_t_2, __pyx_n_s_cl); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 114, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_t_3 = __Pyx_PyObject_GetAttrStr(__pyx_t_2, __pyx_n_s_Context); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 114, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; if (unlikely(__pyx_v_device == Py_None)) { PyErr_SetString(PyExc_TypeError, "'NoneType' object is not subscriptable"); __PYX_ERR(0, 114, __pyx_L1_error) } __pyx_t_2 = __Pyx_PyDict_GetItem(__pyx_v_device, __pyx_n_u_device); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 114, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_t_4 = PyList_New(1); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 114, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_GIVEREF(__pyx_t_2); if (__Pyx_PyList_SET_ITEM(__pyx_t_4, 0, __pyx_t_2)) __PYX_ERR(0, 114, __pyx_L1_error); __pyx_t_2 = 0; __pyx_t_2 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_3))) { __pyx_t_2 = PyMethod_GET_SELF(__pyx_t_3); if (likely(__pyx_t_2)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_3); __Pyx_INCREF(__pyx_t_2); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_3, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[2] = {__pyx_t_2, __pyx_t_4}; __pyx_t_1 = __Pyx_PyObject_FastCall(__pyx_t_3, __pyx_callargs+1-__pyx_t_5, 1+__pyx_t_5); __Pyx_XDECREF(__pyx_t_2); __pyx_t_2 = 0; __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 114, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; } __pyx_v_cl_ctx = __pyx_t_1; __pyx_t_1 = 0;
+115: dc = device['device']
if (unlikely(__pyx_v_device == Py_None)) {
PyErr_SetString(PyExc_TypeError, "'NoneType' object is not subscriptable");
__PYX_ERR(0, 115, __pyx_L1_error)
}
__pyx_t_1 = __Pyx_PyDict_GetItem(__pyx_v_device, __pyx_n_u_device); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 115, __pyx_L1_error)
__Pyx_GOTREF(__pyx_t_1);
__pyx_v_dc = __pyx_t_1;
__pyx_t_1 = 0;
+116: cl_queue = cl.CommandQueue(cl_ctx)
__Pyx_GetModuleGlobalName(__pyx_t_3, __pyx_n_s_cl); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 116, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_4 = __Pyx_PyObject_GetAttrStr(__pyx_t_3, __pyx_n_s_CommandQueue); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 116, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __pyx_t_3 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_4))) { __pyx_t_3 = PyMethod_GET_SELF(__pyx_t_4); if (likely(__pyx_t_3)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_4); __Pyx_INCREF(__pyx_t_3); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_4, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[2] = {__pyx_t_3, __pyx_v_cl_ctx}; __pyx_t_1 = __Pyx_PyObject_FastCall(__pyx_t_4, __pyx_callargs+1-__pyx_t_5, 1+__pyx_t_5); __Pyx_XDECREF(__pyx_t_3); __pyx_t_3 = 0; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 116, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; } __pyx_v_cl_queue = __pyx_t_1; __pyx_t_1 = 0;
117:
118: # Swap row and columns because opencl is strange and stores the
119: # array in a buffer in fortran ordering despite the original
120: # numpy array being in C order.
121: #image = np.ascontiguousarray(np.swapaxes(image, 1, 2), dtype=np.float32)
122:
+123: cdef int nFrames = image.shape[0]
__pyx_v_nFrames = (__pyx_v_image.shape[0]);
+124: cdef int row = image.shape[1]
__pyx_v_row = (__pyx_v_image.shape[1]);
+125: cdef int col = image.shape[2]
__pyx_v_col = (__pyx_v_image.shape[2]);
+126: cdef float [:,:,:] gradient_col = np.zeros_like(image)
__Pyx_GetModuleGlobalName(__pyx_t_4, __pyx_n_s_np); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 126, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __pyx_t_3 = __Pyx_PyObject_GetAttrStr(__pyx_t_4, __pyx_n_s_zeros_like); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 126, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __pyx_t_4 = __pyx_memoryview_fromslice(__pyx_v_image, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 126, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __pyx_t_2 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_3))) { __pyx_t_2 = PyMethod_GET_SELF(__pyx_t_3); if (likely(__pyx_t_2)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_3); __Pyx_INCREF(__pyx_t_2); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_3, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[2] = {__pyx_t_2, __pyx_t_4}; __pyx_t_1 = __Pyx_PyObject_FastCall(__pyx_t_3, __pyx_callargs+1-__pyx_t_5, 1+__pyx_t_5); __Pyx_XDECREF(__pyx_t_2); __pyx_t_2 = 0; __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 126, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; } __pyx_t_6 = __Pyx_PyObject_to_MemoryviewSlice_dsdsds_float(__pyx_t_1, PyBUF_WRITABLE); if (unlikely(!__pyx_t_6.memview)) __PYX_ERR(0, 126, __pyx_L1_error) __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_v_gradient_col = __pyx_t_6; __pyx_t_6.memview = NULL; __pyx_t_6.data = NULL;
+127: cdef float [:,:,:] gradient_row = np.zeros_like(image)
__Pyx_GetModuleGlobalName(__pyx_t_3, __pyx_n_s_np); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 127, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_4 = __Pyx_PyObject_GetAttrStr(__pyx_t_3, __pyx_n_s_zeros_like); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 127, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __pyx_t_3 = __pyx_memoryview_fromslice(__pyx_v_image, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 127, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_2 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_4))) { __pyx_t_2 = PyMethod_GET_SELF(__pyx_t_4); if (likely(__pyx_t_2)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_4); __Pyx_INCREF(__pyx_t_2); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_4, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[2] = {__pyx_t_2, __pyx_t_3}; __pyx_t_1 = __Pyx_PyObject_FastCall(__pyx_t_4, __pyx_callargs+1-__pyx_t_5, 1+__pyx_t_5); __Pyx_XDECREF(__pyx_t_2); __pyx_t_2 = 0; __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 127, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; } __pyx_t_6 = __Pyx_PyObject_to_MemoryviewSlice_dsdsds_float(__pyx_t_1, PyBUF_WRITABLE); if (unlikely(!__pyx_t_6.memview)) __PYX_ERR(0, 127, __pyx_L1_error) __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_v_gradient_row = __pyx_t_6; __pyx_t_6.memview = NULL; __pyx_t_6.data = NULL;
128:
+129: max_slices = int((dc.global_mem_size // (image[0,:,:].nbytes + gradient_col[0,:,:].nbytes + gradient_row[0,:,:].nbytes))/3)
__pyx_t_1 = __Pyx_PyObject_GetAttrStr(__pyx_v_dc, __pyx_n_s_global_mem_size); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 129, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_7.data = __pyx_v_image.data; __pyx_t_7.memview = __pyx_v_image.memview; __PYX_INC_MEMVIEW(&__pyx_t_7, 1); { Py_ssize_t __pyx_tmp_idx = 0; Py_ssize_t __pyx_tmp_stride = __pyx_v_image.strides[0]; __pyx_t_7.data += __pyx_tmp_idx * __pyx_tmp_stride; } __pyx_t_7.shape[0] = __pyx_v_image.shape[1]; __pyx_t_7.strides[0] = __pyx_v_image.strides[1]; __pyx_t_7.suboffsets[0] = -1; __pyx_t_7.shape[1] = __pyx_v_image.shape[2]; __pyx_t_7.strides[1] = __pyx_v_image.strides[2]; __pyx_t_7.suboffsets[1] = -1; __pyx_t_4 = __pyx_memoryview_fromslice(__pyx_t_7, 2, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 129, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __PYX_XCLEAR_MEMVIEW(&__pyx_t_7, 1); __pyx_t_7.memview = NULL; __pyx_t_7.data = NULL; __pyx_t_3 = __Pyx_PyObject_GetAttrStr(__pyx_t_4, __pyx_n_s_nbytes); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 129, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __pyx_t_7.data = __pyx_v_gradient_col.data; __pyx_t_7.memview = __pyx_v_gradient_col.memview; __PYX_INC_MEMVIEW(&__pyx_t_7, 1); { Py_ssize_t __pyx_tmp_idx = 0; Py_ssize_t __pyx_tmp_stride = __pyx_v_gradient_col.strides[0]; __pyx_t_7.data += __pyx_tmp_idx * __pyx_tmp_stride; } __pyx_t_7.shape[0] = __pyx_v_gradient_col.shape[1]; __pyx_t_7.strides[0] = __pyx_v_gradient_col.strides[1]; __pyx_t_7.suboffsets[0] = -1; __pyx_t_7.shape[1] = __pyx_v_gradient_col.shape[2]; __pyx_t_7.strides[1] = __pyx_v_gradient_col.strides[2]; __pyx_t_7.suboffsets[1] = -1; __pyx_t_4 = __pyx_memoryview_fromslice(__pyx_t_7, 2, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 129, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __PYX_XCLEAR_MEMVIEW(&__pyx_t_7, 1); __pyx_t_7.memview = NULL; __pyx_t_7.data = NULL; __pyx_t_2 = __Pyx_PyObject_GetAttrStr(__pyx_t_4, __pyx_n_s_nbytes); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 129, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __pyx_t_4 = PyNumber_Add(__pyx_t_3, __pyx_t_2); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 129, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_7.data = __pyx_v_gradient_row.data; __pyx_t_7.memview = __pyx_v_gradient_row.memview; __PYX_INC_MEMVIEW(&__pyx_t_7, 1); { Py_ssize_t __pyx_tmp_idx = 0; Py_ssize_t __pyx_tmp_stride = __pyx_v_gradient_row.strides[0]; __pyx_t_7.data += __pyx_tmp_idx * __pyx_tmp_stride; } __pyx_t_7.shape[0] = __pyx_v_gradient_row.shape[1]; __pyx_t_7.strides[0] = __pyx_v_gradient_row.strides[1]; __pyx_t_7.suboffsets[0] = -1; __pyx_t_7.shape[1] = __pyx_v_gradient_row.shape[2]; __pyx_t_7.strides[1] = __pyx_v_gradient_row.strides[2]; __pyx_t_7.suboffsets[1] = -1; __pyx_t_2 = __pyx_memoryview_fromslice(__pyx_t_7, 2, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 129, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __PYX_XCLEAR_MEMVIEW(&__pyx_t_7, 1); __pyx_t_7.memview = NULL; __pyx_t_7.data = NULL; __pyx_t_3 = __Pyx_PyObject_GetAttrStr(__pyx_t_2, __pyx_n_s_nbytes); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 129, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = PyNumber_Add(__pyx_t_4, __pyx_t_3); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 129, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __pyx_t_3 = PyNumber_FloorDivide(__pyx_t_1, __pyx_t_2); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 129, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = __Pyx_PyInt_TrueDivideObjC(__pyx_t_3, __pyx_int_3, 3, 0, 0); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 129, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __pyx_t_3 = __Pyx_PyNumber_Int(__pyx_t_2); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 129, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_v_max_slices = __pyx_t_3; __pyx_t_3 = 0;
130: # TODO add exception if max_slices < 1
131:
+132: mf = cl.mem_flags
__Pyx_GetModuleGlobalName(__pyx_t_3, __pyx_n_s_cl); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 132, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_2 = __Pyx_PyObject_GetAttrStr(__pyx_t_3, __pyx_n_s_mem_flags); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 132, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __pyx_v_mf = __pyx_t_2; __pyx_t_2 = 0;
133:
+134: input_opencl = cl.Buffer(cl_ctx, mf.READ_ONLY, image[0:max_slices,:,:].nbytes)
__Pyx_GetModuleGlobalName(__pyx_t_3, __pyx_n_s_cl); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 134, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_1 = __Pyx_PyObject_GetAttrStr(__pyx_t_3, __pyx_n_s_Buffer); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 134, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __pyx_t_3 = __Pyx_PyObject_GetAttrStr(__pyx_v_mf, __pyx_n_s_READ_ONLY); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 134, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_8 = __Pyx_PyIndex_AsSsize_t(__pyx_v_max_slices); if (unlikely((__pyx_t_8 == (Py_ssize_t)-1) && PyErr_Occurred())) __PYX_ERR(0, 134, __pyx_L1_error) __pyx_t_6.data = __pyx_v_image.data; __pyx_t_6.memview = __pyx_v_image.memview; __PYX_INC_MEMVIEW(&__pyx_t_6, 1); __pyx_t_5 = -1; if (unlikely(__pyx_memoryview_slice_memviewslice( &__pyx_t_6, __pyx_v_image.shape[0], __pyx_v_image.strides[0], __pyx_v_image.suboffsets[0], 0, 0, &__pyx_t_5, 0, __pyx_t_8, 0, 1, 1, 0, 1) < 0)) { __PYX_ERR(0, 134, __pyx_L1_error) } __pyx_t_6.shape[1] = __pyx_v_image.shape[1]; __pyx_t_6.strides[1] = __pyx_v_image.strides[1]; __pyx_t_6.suboffsets[1] = -1; __pyx_t_6.shape[2] = __pyx_v_image.shape[2]; __pyx_t_6.strides[2] = __pyx_v_image.strides[2]; __pyx_t_6.suboffsets[2] = -1; __pyx_t_4 = __pyx_memoryview_fromslice(__pyx_t_6, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 134, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __PYX_XCLEAR_MEMVIEW(&__pyx_t_6, 1); __pyx_t_6.memview = NULL; __pyx_t_6.data = NULL; __pyx_t_9 = __Pyx_PyObject_GetAttrStr(__pyx_t_4, __pyx_n_s_nbytes); if (unlikely(!__pyx_t_9)) __PYX_ERR(0, 134, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_9); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __pyx_t_4 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_1))) { __pyx_t_4 = PyMethod_GET_SELF(__pyx_t_1); if (likely(__pyx_t_4)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_1); __Pyx_INCREF(__pyx_t_4); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_1, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[4] = {__pyx_t_4, __pyx_v_cl_ctx, __pyx_t_3, __pyx_t_9}; __pyx_t_2 = __Pyx_PyObject_FastCall(__pyx_t_1, __pyx_callargs+1-__pyx_t_5, 3+__pyx_t_5); __Pyx_XDECREF(__pyx_t_4); __pyx_t_4 = 0; __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __Pyx_DECREF(__pyx_t_9); __pyx_t_9 = 0; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 134, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; } __pyx_v_input_opencl = __pyx_t_2; __pyx_t_2 = 0;
+135: cl.enqueue_copy(cl_queue, input_opencl, image[0:max_slices,:,:]).wait()
__Pyx_GetModuleGlobalName(__pyx_t_9, __pyx_n_s_cl); if (unlikely(!__pyx_t_9)) __PYX_ERR(0, 135, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_9); __pyx_t_3 = __Pyx_PyObject_GetAttrStr(__pyx_t_9, __pyx_n_s_enqueue_copy); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 135, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_DECREF(__pyx_t_9); __pyx_t_9 = 0; __pyx_t_8 = __Pyx_PyIndex_AsSsize_t(__pyx_v_max_slices); if (unlikely((__pyx_t_8 == (Py_ssize_t)-1) && PyErr_Occurred())) __PYX_ERR(0, 135, __pyx_L1_error) __pyx_t_6.data = __pyx_v_image.data; __pyx_t_6.memview = __pyx_v_image.memview; __PYX_INC_MEMVIEW(&__pyx_t_6, 1); __pyx_t_5 = -1; if (unlikely(__pyx_memoryview_slice_memviewslice( &__pyx_t_6, __pyx_v_image.shape[0], __pyx_v_image.strides[0], __pyx_v_image.suboffsets[0], 0, 0, &__pyx_t_5, 0, __pyx_t_8, 0, 1, 1, 0, 1) < 0)) { __PYX_ERR(0, 135, __pyx_L1_error) } __pyx_t_6.shape[1] = __pyx_v_image.shape[1]; __pyx_t_6.strides[1] = __pyx_v_image.strides[1]; __pyx_t_6.suboffsets[1] = -1; __pyx_t_6.shape[2] = __pyx_v_image.shape[2]; __pyx_t_6.strides[2] = __pyx_v_image.strides[2]; __pyx_t_6.suboffsets[2] = -1; __pyx_t_9 = __pyx_memoryview_fromslice(__pyx_t_6, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_9)) __PYX_ERR(0, 135, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_9); __PYX_XCLEAR_MEMVIEW(&__pyx_t_6, 1); __pyx_t_6.memview = NULL; __pyx_t_6.data = NULL; __pyx_t_4 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_3))) { __pyx_t_4 = PyMethod_GET_SELF(__pyx_t_3); if (likely(__pyx_t_4)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_3); __Pyx_INCREF(__pyx_t_4); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_3, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[4] = {__pyx_t_4, __pyx_v_cl_queue, __pyx_v_input_opencl, __pyx_t_9}; __pyx_t_1 = __Pyx_PyObject_FastCall(__pyx_t_3, __pyx_callargs+1-__pyx_t_5, 3+__pyx_t_5); __Pyx_XDECREF(__pyx_t_4); __pyx_t_4 = 0; __Pyx_DECREF(__pyx_t_9); __pyx_t_9 = 0; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 135, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; } __pyx_t_3 = __Pyx_PyObject_GetAttrStr(__pyx_t_1, __pyx_n_s_wait); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 135, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_t_1 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (likely(PyMethod_Check(__pyx_t_3))) { __pyx_t_1 = PyMethod_GET_SELF(__pyx_t_3); if (likely(__pyx_t_1)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_3); __Pyx_INCREF(__pyx_t_1); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_3, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[1] = {__pyx_t_1, }; __pyx_t_2 = __Pyx_PyObject_FastCall(__pyx_t_3, __pyx_callargs+1-__pyx_t_5, 0+__pyx_t_5); __Pyx_XDECREF(__pyx_t_1); __pyx_t_1 = 0; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 135, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; } __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0;
+136: output_opencl_col = cl.Buffer(cl_ctx, mf.WRITE_ONLY, gradient_col[0:max_slices,:,:].nbytes)
__Pyx_GetModuleGlobalName(__pyx_t_3, __pyx_n_s_cl); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 136, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_1 = __Pyx_PyObject_GetAttrStr(__pyx_t_3, __pyx_n_s_Buffer); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 136, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __pyx_t_3 = __Pyx_PyObject_GetAttrStr(__pyx_v_mf, __pyx_n_s_WRITE_ONLY); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 136, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __pyx_t_8 = __Pyx_PyIndex_AsSsize_t(__pyx_v_max_slices); if (unlikely((__pyx_t_8 == (Py_ssize_t)-1) && PyErr_Occurred())) __PYX_ERR(0, 136, __pyx_L1_error) __pyx_t_6.data = __pyx_v_gradient_col.data; __pyx_t_6.memview = __pyx_v_gradient_col.memview; __PYX_INC_MEMVIEW(&__pyx_t_6, 1); __pyx_t_5 = -1; if (unlikely(__pyx_memoryview_slice_memviewslice( &__pyx_t_6, __pyx_v_gradient_col.shape[0], __pyx_v_gradient_col.strides[0], __pyx_v_gradient_col.suboffsets[0], 0, 0, &__pyx_t_5, 0, __pyx_t_8, 0, 1, 1, 0, 1) < 0)) { __PYX_ERR(0, 136, __pyx_L1_error) } __pyx_t_6.shape[1] = __pyx_v_gradient_col.shape[1]; __pyx_t_6.strides[1] = __pyx_v_gradient_col.strides[1]; __pyx_t_6.suboffsets[1] = -1; __pyx_t_6.shape[2] = __pyx_v_gradient_col.shape[2]; __pyx_t_6.strides[2] = __pyx_v_gradient_col.strides[2]; __pyx_t_6.suboffsets[2] = -1; __pyx_t_9 = __pyx_memoryview_fromslice(__pyx_t_6, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_9)) __PYX_ERR(0, 136, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_9); __PYX_XCLEAR_MEMVIEW(&__pyx_t_6, 1); __pyx_t_6.memview = NULL; __pyx_t_6.data = NULL; __pyx_t_4 = __Pyx_PyObject_GetAttrStr(__pyx_t_9, __pyx_n_s_nbytes); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 136, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_DECREF(__pyx_t_9); __pyx_t_9 = 0; __pyx_t_9 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_1))) { __pyx_t_9 = PyMethod_GET_SELF(__pyx_t_1); if (likely(__pyx_t_9)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_1); __Pyx_INCREF(__pyx_t_9); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_1, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[4] = {__pyx_t_9, __pyx_v_cl_ctx, __pyx_t_3, __pyx_t_4}; __pyx_t_2 = __Pyx_PyObject_FastCall(__pyx_t_1, __pyx_callargs+1-__pyx_t_5, 3+__pyx_t_5); __Pyx_XDECREF(__pyx_t_9); __pyx_t_9 = 0; __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 136, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; } __pyx_v_output_opencl_col = __pyx_t_2; __pyx_t_2 = 0;
+137: output_opencl_row = cl.Buffer(cl_ctx, mf.WRITE_ONLY, gradient_row[0:max_slices, :, :].nbytes)
__Pyx_GetModuleGlobalName(__pyx_t_1, __pyx_n_s_cl); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 137, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_4 = __Pyx_PyObject_GetAttrStr(__pyx_t_1, __pyx_n_s_Buffer); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 137, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __pyx_t_1 = __Pyx_PyObject_GetAttrStr(__pyx_v_mf, __pyx_n_s_WRITE_ONLY); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 137, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_8 = __Pyx_PyIndex_AsSsize_t(__pyx_v_max_slices); if (unlikely((__pyx_t_8 == (Py_ssize_t)-1) && PyErr_Occurred())) __PYX_ERR(0, 137, __pyx_L1_error) __pyx_t_6.data = __pyx_v_gradient_row.data; __pyx_t_6.memview = __pyx_v_gradient_row.memview; __PYX_INC_MEMVIEW(&__pyx_t_6, 1); __pyx_t_5 = -1; if (unlikely(__pyx_memoryview_slice_memviewslice( &__pyx_t_6, __pyx_v_gradient_row.shape[0], __pyx_v_gradient_row.strides[0], __pyx_v_gradient_row.suboffsets[0], 0, 0, &__pyx_t_5, 0, __pyx_t_8, 0, 1, 1, 0, 1) < 0)) { __PYX_ERR(0, 137, __pyx_L1_error) } __pyx_t_6.shape[1] = __pyx_v_gradient_row.shape[1]; __pyx_t_6.strides[1] = __pyx_v_gradient_row.strides[1]; __pyx_t_6.suboffsets[1] = -1; __pyx_t_6.shape[2] = __pyx_v_gradient_row.shape[2]; __pyx_t_6.strides[2] = __pyx_v_gradient_row.strides[2]; __pyx_t_6.suboffsets[2] = -1; __pyx_t_3 = __pyx_memoryview_fromslice(__pyx_t_6, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 137, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __PYX_XCLEAR_MEMVIEW(&__pyx_t_6, 1); __pyx_t_6.memview = NULL; __pyx_t_6.data = NULL; __pyx_t_9 = __Pyx_PyObject_GetAttrStr(__pyx_t_3, __pyx_n_s_nbytes); if (unlikely(!__pyx_t_9)) __PYX_ERR(0, 137, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_9); __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __pyx_t_3 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_4))) { __pyx_t_3 = PyMethod_GET_SELF(__pyx_t_4); if (likely(__pyx_t_3)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_4); __Pyx_INCREF(__pyx_t_3); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_4, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[4] = {__pyx_t_3, __pyx_v_cl_ctx, __pyx_t_1, __pyx_t_9}; __pyx_t_2 = __Pyx_PyObject_FastCall(__pyx_t_4, __pyx_callargs+1-__pyx_t_5, 3+__pyx_t_5); __Pyx_XDECREF(__pyx_t_3); __pyx_t_3 = 0; __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; __Pyx_DECREF(__pyx_t_9); __pyx_t_9 = 0; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 137, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; } __pyx_v_output_opencl_row = __pyx_t_2; __pyx_t_2 = 0;
138:
+139: code = self._get_cl_code("_le_roberts_cross_gradients.cl", device['DP'])
__pyx_t_4 = __Pyx_PyObject_GetAttrStr(__pyx_v_self, __pyx_n_s_get_cl_code); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 139, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); if (unlikely(__pyx_v_device == Py_None)) { PyErr_SetString(PyExc_TypeError, "'NoneType' object is not subscriptable"); __PYX_ERR(0, 139, __pyx_L1_error) } __pyx_t_9 = __Pyx_PyDict_GetItem(__pyx_v_device, __pyx_n_u_DP); if (unlikely(!__pyx_t_9)) __PYX_ERR(0, 139, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_9); __pyx_t_1 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (likely(PyMethod_Check(__pyx_t_4))) { __pyx_t_1 = PyMethod_GET_SELF(__pyx_t_4); if (likely(__pyx_t_1)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_4); __Pyx_INCREF(__pyx_t_1); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_4, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[3] = {__pyx_t_1, __pyx_kp_u_le_roberts_cross_gradients_cl, __pyx_t_9}; __pyx_t_2 = __Pyx_PyObject_FastCall(__pyx_t_4, __pyx_callargs+1-__pyx_t_5, 2+__pyx_t_5); __Pyx_XDECREF(__pyx_t_1); __pyx_t_1 = 0; __Pyx_DECREF(__pyx_t_9); __pyx_t_9 = 0; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 139, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; } __pyx_v_code = __pyx_t_2; __pyx_t_2 = 0;
+140: prg = cl.Program(cl_ctx, code).build()
__Pyx_GetModuleGlobalName(__pyx_t_9, __pyx_n_s_cl); if (unlikely(!__pyx_t_9)) __PYX_ERR(0, 140, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_9); __pyx_t_1 = __Pyx_PyObject_GetAttrStr(__pyx_t_9, __pyx_n_s_Program); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 140, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_9); __pyx_t_9 = 0; __pyx_t_9 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_1))) { __pyx_t_9 = PyMethod_GET_SELF(__pyx_t_1); if (likely(__pyx_t_9)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_1); __Pyx_INCREF(__pyx_t_9); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_1, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[3] = {__pyx_t_9, __pyx_v_cl_ctx, __pyx_v_code}; __pyx_t_4 = __Pyx_PyObject_FastCall(__pyx_t_1, __pyx_callargs+1-__pyx_t_5, 2+__pyx_t_5); __Pyx_XDECREF(__pyx_t_9); __pyx_t_9 = 0; if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 140, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; } __pyx_t_1 = __Pyx_PyObject_GetAttrStr(__pyx_t_4, __pyx_n_s_build); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 140, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __pyx_t_4 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (likely(PyMethod_Check(__pyx_t_1))) { __pyx_t_4 = PyMethod_GET_SELF(__pyx_t_1); if (likely(__pyx_t_4)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_1); __Pyx_INCREF(__pyx_t_4); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_1, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[1] = {__pyx_t_4, }; __pyx_t_2 = __Pyx_PyObject_FastCall(__pyx_t_1, __pyx_callargs+1-__pyx_t_5, 0+__pyx_t_5); __Pyx_XDECREF(__pyx_t_4); __pyx_t_4 = 0; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 140, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; } __pyx_v_prg = __pyx_t_2; __pyx_t_2 = 0;
+141: knl = prg.gradient_roberts_cross
__pyx_t_2 = __Pyx_PyObject_GetAttrStr(__pyx_v_prg, __pyx_n_s_gradient_roberts_cross); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 141, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_v_knl = __pyx_t_2; __pyx_t_2 = 0;
142:
+143: for i in range(0, image.shape[0]-1, max_slices):
__pyx_t_2 = PyInt_FromSsize_t(((__pyx_v_image.shape[0]) - 1)); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 143, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_t_1 = PyTuple_New(3); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 143, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_INCREF(__pyx_int_0); __Pyx_GIVEREF(__pyx_int_0); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_1, 0, __pyx_int_0)) __PYX_ERR(0, 143, __pyx_L1_error); __Pyx_GIVEREF(__pyx_t_2); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_1, 1, __pyx_t_2)) __PYX_ERR(0, 143, __pyx_L1_error); __Pyx_INCREF(__pyx_v_max_slices); __Pyx_GIVEREF(__pyx_v_max_slices); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_1, 2, __pyx_v_max_slices)) __PYX_ERR(0, 143, __pyx_L1_error); __pyx_t_2 = 0; __pyx_t_2 = __Pyx_PyObject_Call(__pyx_builtin_range, __pyx_t_1, NULL); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 143, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0; if (likely(PyList_CheckExact(__pyx_t_2)) || PyTuple_CheckExact(__pyx_t_2)) { __pyx_t_1 = __pyx_t_2; __Pyx_INCREF(__pyx_t_1); __pyx_t_8 = 0; __pyx_t_10 = NULL; } else { __pyx_t_8 = -1; __pyx_t_1 = PyObject_GetIter(__pyx_t_2); if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 143, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_10 = __Pyx_PyObject_GetIterNextFunc(__pyx_t_1); if (unlikely(!__pyx_t_10)) __PYX_ERR(0, 143, __pyx_L1_error) } __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; for (;;) { if (likely(!__pyx_t_10)) { if (likely(PyList_CheckExact(__pyx_t_1))) { if (__pyx_t_8 >= PyList_GET_SIZE(__pyx_t_1)) break; #if CYTHON_ASSUME_SAFE_MACROS && !CYTHON_AVOID_BORROWED_REFS __pyx_t_2 = PyList_GET_ITEM(__pyx_t_1, __pyx_t_8); __Pyx_INCREF(__pyx_t_2); __pyx_t_8++; if (unlikely((0 < 0))) __PYX_ERR(0, 143, __pyx_L1_error) #else __pyx_t_2 = PySequence_ITEM(__pyx_t_1, __pyx_t_8); __pyx_t_8++; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 143, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); #endif } else { if (__pyx_t_8 >= PyTuple_GET_SIZE(__pyx_t_1)) break; #if CYTHON_ASSUME_SAFE_MACROS && !CYTHON_AVOID_BORROWED_REFS __pyx_t_2 = PyTuple_GET_ITEM(__pyx_t_1, __pyx_t_8); __Pyx_INCREF(__pyx_t_2); __pyx_t_8++; if (unlikely((0 < 0))) __PYX_ERR(0, 143, __pyx_L1_error) #else __pyx_t_2 = PySequence_ITEM(__pyx_t_1, __pyx_t_8); __pyx_t_8++; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 143, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); #endif } } else { __pyx_t_2 = __pyx_t_10(__pyx_t_1); if (unlikely(!__pyx_t_2)) { PyObject* exc_type = PyErr_Occurred(); if (exc_type) { if (likely(__Pyx_PyErr_GivenExceptionMatches(exc_type, PyExc_StopIteration))) PyErr_Clear(); else __PYX_ERR(0, 143, __pyx_L1_error) } break; } __Pyx_GOTREF(__pyx_t_2); } __Pyx_XDECREF_SET(__pyx_v_i, __pyx_t_2); __pyx_t_2 = 0; /* … */ } __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0;
+144: if image.shape[0] - i >= max_slices:
__pyx_t_2 = PyInt_FromSsize_t((__pyx_v_image.shape[0])); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 144, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_t_4 = PyNumber_Subtract(__pyx_t_2, __pyx_v_i); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 144, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = PyObject_RichCompare(__pyx_t_4, __pyx_v_max_slices, Py_GE); __Pyx_XGOTREF(__pyx_t_2); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 144, __pyx_L1_error) __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __pyx_t_11 = __Pyx_PyObject_IsTrue(__pyx_t_2); if (unlikely((__pyx_t_11 < 0))) __PYX_ERR(0, 144, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; if (__pyx_t_11) { /* … */ goto __pyx_L5; }
+145: n_slices = max_slices
__Pyx_INCREF(__pyx_v_max_slices); __Pyx_XDECREF_SET(__pyx_v_n_slices, __pyx_v_max_slices);
146: else:
+147: n_slices = image.shape[0] - i
/*else*/ {
__pyx_t_2 = PyInt_FromSsize_t((__pyx_v_image.shape[0])); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 147, __pyx_L1_error)
__Pyx_GOTREF(__pyx_t_2);
__pyx_t_4 = PyNumber_Subtract(__pyx_t_2, __pyx_v_i); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 147, __pyx_L1_error)
__Pyx_GOTREF(__pyx_t_4);
__Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0;
__Pyx_XDECREF_SET(__pyx_v_n_slices, __pyx_t_4);
__pyx_t_4 = 0;
}
__pyx_L5:;
148: knl(cl_queue,
+149: (n_slices,),
__pyx_t_9 = PyTuple_New(1); if (unlikely(!__pyx_t_9)) __PYX_ERR(0, 149, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_9); __Pyx_INCREF(__pyx_v_n_slices); __Pyx_GIVEREF(__pyx_v_n_slices); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_9, 0, __pyx_v_n_slices)) __PYX_ERR(0, 149, __pyx_L1_error);
150: None,
151: input_opencl,
152: output_opencl_col,
153: output_opencl_row,
+154: np.int32(row),
__Pyx_GetModuleGlobalName(__pyx_t_12, __pyx_n_s_np); if (unlikely(!__pyx_t_12)) __PYX_ERR(0, 154, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_12); __pyx_t_13 = __Pyx_PyObject_GetAttrStr(__pyx_t_12, __pyx_n_s_int32); if (unlikely(!__pyx_t_13)) __PYX_ERR(0, 154, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_13); __Pyx_DECREF(__pyx_t_12); __pyx_t_12 = 0; __pyx_t_12 = __Pyx_PyInt_From_int(__pyx_v_row); if (unlikely(!__pyx_t_12)) __PYX_ERR(0, 154, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_12); __pyx_t_14 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_13))) { __pyx_t_14 = PyMethod_GET_SELF(__pyx_t_13); if (likely(__pyx_t_14)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_13); __Pyx_INCREF(__pyx_t_14); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_13, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[2] = {__pyx_t_14, __pyx_t_12}; __pyx_t_3 = __Pyx_PyObject_FastCall(__pyx_t_13, __pyx_callargs+1-__pyx_t_5, 1+__pyx_t_5); __Pyx_XDECREF(__pyx_t_14); __pyx_t_14 = 0; __Pyx_DECREF(__pyx_t_12); __pyx_t_12 = 0; if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 154, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_DECREF(__pyx_t_13); __pyx_t_13 = 0; }
+155: np.int32(col)).wait()
__Pyx_GetModuleGlobalName(__pyx_t_12, __pyx_n_s_np); if (unlikely(!__pyx_t_12)) __PYX_ERR(0, 155, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_12); __pyx_t_14 = __Pyx_PyObject_GetAttrStr(__pyx_t_12, __pyx_n_s_int32); if (unlikely(!__pyx_t_14)) __PYX_ERR(0, 155, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_14); __Pyx_DECREF(__pyx_t_12); __pyx_t_12 = 0; __pyx_t_12 = __Pyx_PyInt_From_int(__pyx_v_col); if (unlikely(!__pyx_t_12)) __PYX_ERR(0, 155, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_12); __pyx_t_15 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_14))) { __pyx_t_15 = PyMethod_GET_SELF(__pyx_t_14); if (likely(__pyx_t_15)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_14); __Pyx_INCREF(__pyx_t_15); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_14, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[2] = {__pyx_t_15, __pyx_t_12}; __pyx_t_13 = __Pyx_PyObject_FastCall(__pyx_t_14, __pyx_callargs+1-__pyx_t_5, 1+__pyx_t_5); __Pyx_XDECREF(__pyx_t_15); __pyx_t_15 = 0; __Pyx_DECREF(__pyx_t_12); __pyx_t_12 = 0; if (unlikely(!__pyx_t_13)) __PYX_ERR(0, 155, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_13); __Pyx_DECREF(__pyx_t_14); __pyx_t_14 = 0; } __Pyx_INCREF(__pyx_v_knl); __pyx_t_14 = __pyx_v_knl; __pyx_t_12 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (likely(PyMethod_Check(__pyx_t_14))) { __pyx_t_12 = PyMethod_GET_SELF(__pyx_t_14); if (likely(__pyx_t_12)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_14); __Pyx_INCREF(__pyx_t_12); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_14, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[9] = {__pyx_t_12, __pyx_v_cl_queue, __pyx_t_9, Py_None, __pyx_v_input_opencl, __pyx_v_output_opencl_col, __pyx_v_output_opencl_row, __pyx_t_3, __pyx_t_13}; __pyx_t_2 = __Pyx_PyObject_FastCall(__pyx_t_14, __pyx_callargs+1-__pyx_t_5, 8+__pyx_t_5); __Pyx_XDECREF(__pyx_t_12); __pyx_t_12 = 0; __Pyx_DECREF(__pyx_t_9); __pyx_t_9 = 0; __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __Pyx_DECREF(__pyx_t_13); __pyx_t_13 = 0; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 148, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_14); __pyx_t_14 = 0; } __pyx_t_14 = __Pyx_PyObject_GetAttrStr(__pyx_t_2, __pyx_n_s_wait); if (unlikely(!__pyx_t_14)) __PYX_ERR(0, 155, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_14); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (likely(PyMethod_Check(__pyx_t_14))) { __pyx_t_2 = PyMethod_GET_SELF(__pyx_t_14); if (likely(__pyx_t_2)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_14); __Pyx_INCREF(__pyx_t_2); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_14, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[1] = {__pyx_t_2, }; __pyx_t_4 = __Pyx_PyObject_FastCall(__pyx_t_14, __pyx_callargs+1-__pyx_t_5, 0+__pyx_t_5); __Pyx_XDECREF(__pyx_t_2); __pyx_t_2 = 0; if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 155, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_DECREF(__pyx_t_14); __pyx_t_14 = 0; } __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0;
156:
+157: cl.enqueue_copy(cl_queue, gradient_col[i:i+n_slices,:,:], output_opencl_col).wait()
__Pyx_GetModuleGlobalName(__pyx_t_2, __pyx_n_s_cl); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 157, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_t_13 = __Pyx_PyObject_GetAttrStr(__pyx_t_2, __pyx_n_s_enqueue_copy); if (unlikely(!__pyx_t_13)) __PYX_ERR(0, 157, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_13); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_16 = __Pyx_PyIndex_AsSsize_t(__pyx_v_i); if (unlikely((__pyx_t_16 == (Py_ssize_t)-1) && PyErr_Occurred())) __PYX_ERR(0, 157, __pyx_L1_error) __pyx_t_2 = PyNumber_Add(__pyx_v_i, __pyx_v_n_slices); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 157, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __pyx_t_17 = __Pyx_PyIndex_AsSsize_t(__pyx_t_2); if (unlikely((__pyx_t_17 == (Py_ssize_t)-1) && PyErr_Occurred())) __PYX_ERR(0, 157, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_6.data = __pyx_v_gradient_col.data; __pyx_t_6.memview = __pyx_v_gradient_col.memview; __PYX_INC_MEMVIEW(&__pyx_t_6, 1); __pyx_t_5 = -1; if (unlikely(__pyx_memoryview_slice_memviewslice( &__pyx_t_6, __pyx_v_gradient_col.shape[0], __pyx_v_gradient_col.strides[0], __pyx_v_gradient_col.suboffsets[0], 0, 0, &__pyx_t_5, __pyx_t_16, __pyx_t_17, 0, 1, 1, 0, 1) < 0)) { __PYX_ERR(0, 157, __pyx_L1_error) } __pyx_t_6.shape[1] = __pyx_v_gradient_col.shape[1]; __pyx_t_6.strides[1] = __pyx_v_gradient_col.strides[1]; __pyx_t_6.suboffsets[1] = -1; __pyx_t_6.shape[2] = __pyx_v_gradient_col.shape[2]; __pyx_t_6.strides[2] = __pyx_v_gradient_col.strides[2]; __pyx_t_6.suboffsets[2] = -1; __pyx_t_2 = __pyx_memoryview_fromslice(__pyx_t_6, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 157, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __PYX_XCLEAR_MEMVIEW(&__pyx_t_6, 1); __pyx_t_6.memview = NULL; __pyx_t_6.data = NULL; __pyx_t_3 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_13))) { __pyx_t_3 = PyMethod_GET_SELF(__pyx_t_13); if (likely(__pyx_t_3)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_13); __Pyx_INCREF(__pyx_t_3); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_13, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[4] = {__pyx_t_3, __pyx_v_cl_queue, __pyx_t_2, __pyx_v_output_opencl_col}; __pyx_t_14 = __Pyx_PyObject_FastCall(__pyx_t_13, __pyx_callargs+1-__pyx_t_5, 3+__pyx_t_5); __Pyx_XDECREF(__pyx_t_3); __pyx_t_3 = 0; __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; if (unlikely(!__pyx_t_14)) __PYX_ERR(0, 157, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_14); __Pyx_DECREF(__pyx_t_13); __pyx_t_13 = 0; } __pyx_t_13 = __Pyx_PyObject_GetAttrStr(__pyx_t_14, __pyx_n_s_wait); if (unlikely(!__pyx_t_13)) __PYX_ERR(0, 157, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_13); __Pyx_DECREF(__pyx_t_14); __pyx_t_14 = 0; __pyx_t_14 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (likely(PyMethod_Check(__pyx_t_13))) { __pyx_t_14 = PyMethod_GET_SELF(__pyx_t_13); if (likely(__pyx_t_14)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_13); __Pyx_INCREF(__pyx_t_14); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_13, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[1] = {__pyx_t_14, }; __pyx_t_4 = __Pyx_PyObject_FastCall(__pyx_t_13, __pyx_callargs+1-__pyx_t_5, 0+__pyx_t_5); __Pyx_XDECREF(__pyx_t_14); __pyx_t_14 = 0; if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 157, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_DECREF(__pyx_t_13); __pyx_t_13 = 0; } __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0;
+158: cl.enqueue_copy(cl_queue, gradient_row[i:i+n_slices,:,:], output_opencl_row).wait()
__Pyx_GetModuleGlobalName(__pyx_t_14, __pyx_n_s_cl); if (unlikely(!__pyx_t_14)) __PYX_ERR(0, 158, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_14); __pyx_t_2 = __Pyx_PyObject_GetAttrStr(__pyx_t_14, __pyx_n_s_enqueue_copy); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 158, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_14); __pyx_t_14 = 0; __pyx_t_17 = __Pyx_PyIndex_AsSsize_t(__pyx_v_i); if (unlikely((__pyx_t_17 == (Py_ssize_t)-1) && PyErr_Occurred())) __PYX_ERR(0, 158, __pyx_L1_error) __pyx_t_14 = PyNumber_Add(__pyx_v_i, __pyx_v_n_slices); if (unlikely(!__pyx_t_14)) __PYX_ERR(0, 158, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_14); __pyx_t_16 = __Pyx_PyIndex_AsSsize_t(__pyx_t_14); if (unlikely((__pyx_t_16 == (Py_ssize_t)-1) && PyErr_Occurred())) __PYX_ERR(0, 158, __pyx_L1_error) __Pyx_DECREF(__pyx_t_14); __pyx_t_14 = 0; __pyx_t_6.data = __pyx_v_gradient_row.data; __pyx_t_6.memview = __pyx_v_gradient_row.memview; __PYX_INC_MEMVIEW(&__pyx_t_6, 1); __pyx_t_5 = -1; if (unlikely(__pyx_memoryview_slice_memviewslice( &__pyx_t_6, __pyx_v_gradient_row.shape[0], __pyx_v_gradient_row.strides[0], __pyx_v_gradient_row.suboffsets[0], 0, 0, &__pyx_t_5, __pyx_t_17, __pyx_t_16, 0, 1, 1, 0, 1) < 0)) { __PYX_ERR(0, 158, __pyx_L1_error) } __pyx_t_6.shape[1] = __pyx_v_gradient_row.shape[1]; __pyx_t_6.strides[1] = __pyx_v_gradient_row.strides[1]; __pyx_t_6.suboffsets[1] = -1; __pyx_t_6.shape[2] = __pyx_v_gradient_row.shape[2]; __pyx_t_6.strides[2] = __pyx_v_gradient_row.strides[2]; __pyx_t_6.suboffsets[2] = -1; __pyx_t_14 = __pyx_memoryview_fromslice(__pyx_t_6, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_14)) __PYX_ERR(0, 158, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_14); __PYX_XCLEAR_MEMVIEW(&__pyx_t_6, 1); __pyx_t_6.memview = NULL; __pyx_t_6.data = NULL; __pyx_t_3 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_2))) { __pyx_t_3 = PyMethod_GET_SELF(__pyx_t_2); if (likely(__pyx_t_3)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_2); __Pyx_INCREF(__pyx_t_3); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_2, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[4] = {__pyx_t_3, __pyx_v_cl_queue, __pyx_t_14, __pyx_v_output_opencl_row}; __pyx_t_13 = __Pyx_PyObject_FastCall(__pyx_t_2, __pyx_callargs+1-__pyx_t_5, 3+__pyx_t_5); __Pyx_XDECREF(__pyx_t_3); __pyx_t_3 = 0; __Pyx_DECREF(__pyx_t_14); __pyx_t_14 = 0; if (unlikely(!__pyx_t_13)) __PYX_ERR(0, 158, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_13); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; } __pyx_t_2 = __Pyx_PyObject_GetAttrStr(__pyx_t_13, __pyx_n_s_wait); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 158, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_13); __pyx_t_13 = 0; __pyx_t_13 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (likely(PyMethod_Check(__pyx_t_2))) { __pyx_t_13 = PyMethod_GET_SELF(__pyx_t_2); if (likely(__pyx_t_13)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_2); __Pyx_INCREF(__pyx_t_13); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_2, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[1] = {__pyx_t_13, }; __pyx_t_4 = __Pyx_PyObject_FastCall(__pyx_t_2, __pyx_callargs+1-__pyx_t_5, 0+__pyx_t_5); __Pyx_XDECREF(__pyx_t_13); __pyx_t_13 = 0; if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 158, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; } __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0;
+159: if i<=image.shape[0]-max_slices:
__pyx_t_4 = PyInt_FromSsize_t((__pyx_v_image.shape[0])); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 159, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __pyx_t_2 = PyNumber_Subtract(__pyx_t_4, __pyx_v_max_slices); if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 159, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; __pyx_t_4 = PyObject_RichCompare(__pyx_v_i, __pyx_t_2, Py_LE); __Pyx_XGOTREF(__pyx_t_4); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 159, __pyx_L1_error) __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_11 = __Pyx_PyObject_IsTrue(__pyx_t_4); if (unlikely((__pyx_t_11 < 0))) __PYX_ERR(0, 159, __pyx_L1_error) __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; if (__pyx_t_11) { /* … */ }
+160: cl.enqueue_copy(cl_queue, input_opencl, image[i+n_slices:i+2*n_slices,:,:]).wait()
__Pyx_GetModuleGlobalName(__pyx_t_13, __pyx_n_s_cl); if (unlikely(!__pyx_t_13)) __PYX_ERR(0, 160, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_13); __pyx_t_14 = __Pyx_PyObject_GetAttrStr(__pyx_t_13, __pyx_n_s_enqueue_copy); if (unlikely(!__pyx_t_14)) __PYX_ERR(0, 160, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_14); __Pyx_DECREF(__pyx_t_13); __pyx_t_13 = 0; __pyx_t_13 = PyNumber_Add(__pyx_v_i, __pyx_v_n_slices); if (unlikely(!__pyx_t_13)) __PYX_ERR(0, 160, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_13); __pyx_t_16 = __Pyx_PyIndex_AsSsize_t(__pyx_t_13); if (unlikely((__pyx_t_16 == (Py_ssize_t)-1) && PyErr_Occurred())) __PYX_ERR(0, 160, __pyx_L1_error) __Pyx_DECREF(__pyx_t_13); __pyx_t_13 = 0; __pyx_t_13 = __Pyx_PyInt_MultiplyCObj(__pyx_int_2, __pyx_v_n_slices, 2, 0, 0); if (unlikely(!__pyx_t_13)) __PYX_ERR(0, 160, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_13); __pyx_t_3 = PyNumber_Add(__pyx_v_i, __pyx_t_13); if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 160, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __Pyx_DECREF(__pyx_t_13); __pyx_t_13 = 0; __pyx_t_17 = __Pyx_PyIndex_AsSsize_t(__pyx_t_3); if (unlikely((__pyx_t_17 == (Py_ssize_t)-1) && PyErr_Occurred())) __PYX_ERR(0, 160, __pyx_L1_error) __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; __pyx_t_6.data = __pyx_v_image.data; __pyx_t_6.memview = __pyx_v_image.memview; __PYX_INC_MEMVIEW(&__pyx_t_6, 1); __pyx_t_5 = -1; if (unlikely(__pyx_memoryview_slice_memviewslice( &__pyx_t_6, __pyx_v_image.shape[0], __pyx_v_image.strides[0], __pyx_v_image.suboffsets[0], 0, 0, &__pyx_t_5, __pyx_t_16, __pyx_t_17, 0, 1, 1, 0, 1) < 0)) { __PYX_ERR(0, 160, __pyx_L1_error) } __pyx_t_6.shape[1] = __pyx_v_image.shape[1]; __pyx_t_6.strides[1] = __pyx_v_image.strides[1]; __pyx_t_6.suboffsets[1] = -1; __pyx_t_6.shape[2] = __pyx_v_image.shape[2]; __pyx_t_6.strides[2] = __pyx_v_image.strides[2]; __pyx_t_6.suboffsets[2] = -1; __pyx_t_3 = __pyx_memoryview_fromslice(__pyx_t_6, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_3)) __PYX_ERR(0, 160, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_3); __PYX_XCLEAR_MEMVIEW(&__pyx_t_6, 1); __pyx_t_6.memview = NULL; __pyx_t_6.data = NULL; __pyx_t_13 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (unlikely(PyMethod_Check(__pyx_t_14))) { __pyx_t_13 = PyMethod_GET_SELF(__pyx_t_14); if (likely(__pyx_t_13)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_14); __Pyx_INCREF(__pyx_t_13); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_14, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[4] = {__pyx_t_13, __pyx_v_cl_queue, __pyx_v_input_opencl, __pyx_t_3}; __pyx_t_2 = __Pyx_PyObject_FastCall(__pyx_t_14, __pyx_callargs+1-__pyx_t_5, 3+__pyx_t_5); __Pyx_XDECREF(__pyx_t_13); __pyx_t_13 = 0; __Pyx_DECREF(__pyx_t_3); __pyx_t_3 = 0; if (unlikely(!__pyx_t_2)) __PYX_ERR(0, 160, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_2); __Pyx_DECREF(__pyx_t_14); __pyx_t_14 = 0; } __pyx_t_14 = __Pyx_PyObject_GetAttrStr(__pyx_t_2, __pyx_n_s_wait); if (unlikely(!__pyx_t_14)) __PYX_ERR(0, 160, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_14); __Pyx_DECREF(__pyx_t_2); __pyx_t_2 = 0; __pyx_t_2 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (likely(PyMethod_Check(__pyx_t_14))) { __pyx_t_2 = PyMethod_GET_SELF(__pyx_t_14); if (likely(__pyx_t_2)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_14); __Pyx_INCREF(__pyx_t_2); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_14, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[1] = {__pyx_t_2, }; __pyx_t_4 = __Pyx_PyObject_FastCall(__pyx_t_14, __pyx_callargs+1-__pyx_t_5, 0+__pyx_t_5); __Pyx_XDECREF(__pyx_t_2); __pyx_t_2 = 0; if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 160, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_DECREF(__pyx_t_14); __pyx_t_14 = 0; } __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0;
161:
+162: cl_queue.finish()
__pyx_t_14 = __Pyx_PyObject_GetAttrStr(__pyx_v_cl_queue, __pyx_n_s_finish); if (unlikely(!__pyx_t_14)) __PYX_ERR(0, 162, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_14); __pyx_t_2 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (likely(PyMethod_Check(__pyx_t_14))) { __pyx_t_2 = PyMethod_GET_SELF(__pyx_t_14); if (likely(__pyx_t_2)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_14); __Pyx_INCREF(__pyx_t_2); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_14, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[1] = {__pyx_t_2, }; __pyx_t_4 = __Pyx_PyObject_FastCall(__pyx_t_14, __pyx_callargs+1-__pyx_t_5, 0+__pyx_t_5); __Pyx_XDECREF(__pyx_t_2); __pyx_t_2 = 0; if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 162, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __Pyx_DECREF(__pyx_t_14); __pyx_t_14 = 0; } __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0;
163:
+164: input_opencl.release()
__pyx_t_4 = __Pyx_PyObject_GetAttrStr(__pyx_v_input_opencl, __pyx_n_s_release); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 164, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __pyx_t_14 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (likely(PyMethod_Check(__pyx_t_4))) { __pyx_t_14 = PyMethod_GET_SELF(__pyx_t_4); if (likely(__pyx_t_14)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_4); __Pyx_INCREF(__pyx_t_14); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_4, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[1] = {__pyx_t_14, }; __pyx_t_1 = __Pyx_PyObject_FastCall(__pyx_t_4, __pyx_callargs+1-__pyx_t_5, 0+__pyx_t_5); __Pyx_XDECREF(__pyx_t_14); __pyx_t_14 = 0; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 164, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; } __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0;
+165: output_opencl_col.release()
__pyx_t_4 = __Pyx_PyObject_GetAttrStr(__pyx_v_output_opencl_col, __pyx_n_s_release); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 165, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __pyx_t_14 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (likely(PyMethod_Check(__pyx_t_4))) { __pyx_t_14 = PyMethod_GET_SELF(__pyx_t_4); if (likely(__pyx_t_14)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_4); __Pyx_INCREF(__pyx_t_14); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_4, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[1] = {__pyx_t_14, }; __pyx_t_1 = __Pyx_PyObject_FastCall(__pyx_t_4, __pyx_callargs+1-__pyx_t_5, 0+__pyx_t_5); __Pyx_XDECREF(__pyx_t_14); __pyx_t_14 = 0; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 165, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; } __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0;
+166: output_opencl_row.release()
__pyx_t_4 = __Pyx_PyObject_GetAttrStr(__pyx_v_output_opencl_row, __pyx_n_s_release); if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 166, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __pyx_t_14 = NULL; __pyx_t_5 = 0; #if CYTHON_UNPACK_METHODS if (likely(PyMethod_Check(__pyx_t_4))) { __pyx_t_14 = PyMethod_GET_SELF(__pyx_t_4); if (likely(__pyx_t_14)) { PyObject* function = PyMethod_GET_FUNCTION(__pyx_t_4); __Pyx_INCREF(__pyx_t_14); __Pyx_INCREF(function); __Pyx_DECREF_SET(__pyx_t_4, function); __pyx_t_5 = 1; } } #endif { PyObject *__pyx_callargs[1] = {__pyx_t_14, }; __pyx_t_1 = __Pyx_PyObject_FastCall(__pyx_t_4, __pyx_callargs+1-__pyx_t_5, 0+__pyx_t_5); __Pyx_XDECREF(__pyx_t_14); __pyx_t_14 = 0; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 166, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __Pyx_DECREF(__pyx_t_4); __pyx_t_4 = 0; } __Pyx_DECREF(__pyx_t_1); __pyx_t_1 = 0;
167: # Swap rows and columns back
168: #return np.ascontiguousarray(np.swapaxes(gradient_col, 1, 2), dtype=np.float32), np.ascontiguousarray(np.swapaxes(gradient_row, 1, 2), dtype=np.float32)
+169: return gradient_col, gradient_row
__Pyx_XDECREF(__pyx_r); __pyx_t_1 = __pyx_memoryview_fromslice(__pyx_v_gradient_col, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_1)) __PYX_ERR(0, 169, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_1); __pyx_t_4 = __pyx_memoryview_fromslice(__pyx_v_gradient_row, 3, (PyObject *(*)(char *)) __pyx_memview_get_float, (int (*)(char *, PyObject *)) __pyx_memview_set_float, 0);; if (unlikely(!__pyx_t_4)) __PYX_ERR(0, 169, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_4); __pyx_t_14 = PyTuple_New(2); if (unlikely(!__pyx_t_14)) __PYX_ERR(0, 169, __pyx_L1_error) __Pyx_GOTREF(__pyx_t_14); __Pyx_GIVEREF(__pyx_t_1); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_14, 0, __pyx_t_1)) __PYX_ERR(0, 169, __pyx_L1_error); __Pyx_GIVEREF(__pyx_t_4); if (__Pyx_PyTuple_SET_ITEM(__pyx_t_14, 1, __pyx_t_4)) __PYX_ERR(0, 169, __pyx_L1_error); __pyx_t_1 = 0; __pyx_t_4 = 0; __pyx_r = __pyx_t_14; __pyx_t_14 = 0; goto __pyx_L0;